1
Università degli Studi di Milano Bicocca
Facoltà di Scienze Matematiche Fisiche e Naturali
Corso di Laurea in Informatica
Sviluppo di un framework in
OpenCL e C++ per Image
Processing
Supervisori:
Dott. Prof. Gianluigi Ciocca
Dott. Prof. Alessandro Colombo
Relazione della prova
nale di:
Paolo Surricchio
Matricola: 708622
Anno Accademico 2009/2010
2
Is a man not entitled to the sweat of his own brow?
'No!' says the man in Washington, 'It belongs to the poor.'
'No!' says the man in the Vatican, 'It belongs to God.'
'No!' says the man in Moscow, 'It belongs to everyone.'
I rejected those answers; instead, I chose something dierent.
I chose the impossible. ... 1
(Andrew Ryan, a character from Bioshock the videogame )
1 http://www.2kgames.com/bioshock/
3
0.1 Ringraziamenti
Ho voluto dedicare una pagina di ringraziamenti in quanto questa tesi è il risultato del lavoro congiunto di così tante persone, vicine e lontane, conosciuti e
sconosciuti, che sarebbe sbagliato non ringraziare.
Sicuramente il primo ringraziamento va ai Dott.ri Gianluigi Ciocca e Alessandro Colombo che mi hanno dato l'opportunità di lavorare su questo argomento
veramente molto aascinante e, insieme, abbiamo deciso come sviluppare il lavoro, giorno per giorno. Grazie per avermi accompagnato in questa ricerca che
mi ha dato molto più di quanto può sembrare. Avete svolto un ruolo ben oltre
la normale didattica dimostrando una professionalità che solo la passione per
questa scienza può veramente giusticare.
Grazie anche a tutto il team del
laboratorio IVL che condivide con loro professionalità e passione.
Un ringraziamento speciale va a David Gohara, ricercatore al centro di biologia computazionale di St. Louis[13]. Grazie per aver rilasciato gratuitamente
i tutorial di OpenCL, l'unico barlume di luce che ho trovato all'inizio quando
tutto mi sembrava complicato ed oscuro. Ringrazio direttamente tutto lo sta
del sito
http://www.macresearch.org/
che rilasciano gratuitamente strumen-
ti professionali per lavorare in ambito scientico; è grazie alle persone come
loro che l'essere umano continua a crescere ed imparare senza limiti economici,
politici e/o religiosi.
Un altro ringraziamento va al materiale rilasciato da AMD e NVIDIA relativo ad OpenCL. In particolare, ringrazio Justin Hensley, Senior Member dello
sta tecnico per i tutorial ed il materiale, compreso il SDK installato su linux
per lavorare in OpenCL. Un piccolo ringraziamento all'utente del forum nou
per aver integrato il SDK di ATI/AMD di OpenCL in dei pacchetti .deb che
sono comodamente installabili senza dover impostare nient'altro.
Un ultimo ringraziamento, prima di concludere, va a tutti quelli che, dietro
al sipario , ogni giorno lavorano per persone che non vedranno mai. Grazie a
tutti coloro che hanno reso possibile la realizzazione di questa tesina e non sono
stati menzionati.
E per ultimo, ma non per importanza, un ringraziamento speciale alla mia
famiglia. Sapete benissimo cosa avete fatto per me e, alla ne, non esisteranno
parole che possano essere scritte su un foglio di carta per ringraziarvi. In particolare grazie per avere creduto in me, grazie per aver creduto sin dall'inizio che
ce l'avrei fatta, quando neanche io ci credevo, quando neanche io credevo che
fosse possibile.
Grazie per avermi fatto capire il signicato della parola credere .
Grazie a tutti.
Paolo Surricchio.
Indice
0.1
Ringraziamenti
. . . . . . . . . . . . . . . . . . . . . . . . . . . .
3
0.2
Introduzione . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
6
0.2.1
7
Struttura del documento . . . . . . . . . . . . . . . . . . .
1 Architettura delle GPU
1.1
1.2
. . . . . . . . . . . . .
9
1.1.1
Evoluzione delle GPU
. . . . . . . . . . . . . . . . . . . .
9
1.1.2
La situazione odierna
. . . . . . . . . . . . . . . . . . . .
10
Architettura generica di una GPU
1.2.1
1.2.2
1.3
9
La storia delle GPU: radici ed evoluzione
Architettura
. . . . . . . . . . . . . . . . .
11
. . . . . . . . . . . . . . . . . . . . . . . . .
12
1.2.1.1
GPU: Graphics Processing Unit
1.2.1.2
Video Bios
. . . . . . . . . . . . . . . . . . . . .
. . . . . . . . .
13
1.2.1.3
Video Memory . . . . . . . . . . . . . . . . . . .
13
1.2.1.4
RAMDAC
. . . . . . . . . . . . . . . . . . . . .
13
1.2.1.5
Output
. . . . . . . . . . . . . . . . . . . . . . .
14
1.2.1.6
Motherboard Interface . . . . . . . . . . . . . . .
14
1.2.1.7
Cooling device . . . . . . . . . . . . . . . . . . .
14
1.2.1.8
Power Demand . . . . . . . . . . . . . . . . . . .
14
1.2.1.9
Osservazioni
14
. . . . . . . . . . . . . . . . . . . .
13
Pipeline di rendering[11, 12] . . . . . . . . . . . . . . . . .
15
1.2.2.1
Trasformation
15
1.2.2.2
Per-Vertex Lightning
1.2.2.3
Viewing Trasformation
. . . . . . . . . . . . . .
16
1.2.2.4
Primitives Generation . . . . . . . . . . . . . . .
16
1.2.2.5
Projection Trasformation . . . . . . . . . . . . .
16
1.2.2.6
Clipping
. . . . . . . . . . . . . . . . . . . . . .
16
1.2.2.7
Scan Conversion, Rasterization . . . . . . . . . .
17
1.2.2.8
Texturing, fragment shading
. . . . . . . . . . .
17
1.2.2.9
Display . . . . . . . . . . . . . . . . . . . . . . .
17
. . . . . . . . . . . . . . . . . . .
. . . . . . . . . . . . . . .
16
Un esame di una scheda: NVIDIA GeForce GTX 285 . . . . . . .
17
1.3.1
Descrizione generale
. . . . . . . . . . . . . . . . . . . . .
18
1.3.2
Dettagli dell'analisi . . . . . . . . . . . . . . . . . . . . . .
18
1.3.3
Dierenze qualitative fra GPU e CPU
. . . . . . . . . . .
19
Conclusioni dell'analisi fra CPU e GPU . . . . .
22
Il chip GT200: organizzazione logica . . . . . . . . . . . .
22
1.3.3.1
1.3.4
4
5
INDICE
1.3.4.1
1.4
Streaming Multiprocessor . . . . . . . . . . . . .
Conclusioni sull'architettura delle GPU
. . . . . . . . . . . . . .
2 OpenCL: Open Computing Language
2.1
2.2
28
2.1.1
Storia
. . . . . . . . . . . . . . . . . . . . . . . . . . . . .
28
2.1.2
Cos'è
. . . . . . . . . . . . . . . . . . . . . . . . . . . . .
29
2.1.3
Perchè . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
29
2.1.4
Come
30
2.1.5
Perchè OpenCL per Image Processing
. . . . . . . . . . . . . . . . . . . . . . . . . . . . .
. . . . . . . . . . .
31
2.2.1
The OpenCL Architecture . . . . . . . . . . . . . . . . . .
31
2.2.1.1
Modello di piattaforma
31
2.2.1.2
Modello di esecuzione . . . . . . . . . . . . . . .
33
2.2.1.3
Modello di memoria . . . . . . . . . . . . . . . .
34
2.2.1.4
Modello di programmazione . . . . . . . . . . . .
35
. . . . . . . . . . . . . .
2.2.2
Il Framework OpenCL . . . . . . . . . . . . . . . . . . . .
35
2.2.3
Il livello piattaforma di OpenCL
36
. . . . . . . . . . . . . .
2.2.3.1
Richieste sulla piattaforma
. . . . . . . . . . . .
36
2.2.3.2
OpenCL Device
. . . . . . . . . . . . . . . . . .
36
2.2.3.3
Il Contesto . . . . . . . . . . . . . . . . . . . . .
38
The OpenCL Runtime . . . . . . . . . . . . . . . . . . . .
38
2.2.4.1
Code di comando
38
2.2.4.2
Oggetti di memoria
. . . . . . . . . . . . . . . .
38
2.2.4.3
Oggetti programma
. . . . . . . . . . . . . . . .
40
2.2.4.4
Oggetti kernel
. . . . . . . . . . . . . . . . . . .
40
2.2.4.5
Esecuzione dei kernels . . . . . . . . . . . . . . .
40
2.2.4.6
Oggetti evento . . . . . . . . . . . . . . . . . . .
41
2.2.4.7
Flush and nish
. . . . . . . . . . . . . . . . .
. . . . . . . . . . . . . . . . . .
41
Conclusioni . . . . . . . . . . . . . . . . . . . . . . . . . .
42
Un esempio di OpenCL
. . . . . . . . . . . . . . . . . . . . . . .
3.2
42
2.3.1
Il kernel . . . . . . . . . . . . . . . . . . . . . . . . . . . .
42
2.3.2
Inizializzazione delle risorse . . . . . . . . . . . . . . . . .
44
2.3.3
Compilazione . . . . . . . . . . . . . . . . . . . . . . . . .
47
2.3.4
Creazione degli oggetti memoria
47
2.3.5
Esecuzione dei kernel . . . . . . . . . . . . . . . . . . . . .
49
2.3.6
Lettura dei dati e Release delle risorse . . . . . . . . . . .
49
Conclusioni
. . . . . . . . . . . . . .
. . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
3 Progettazione del Framework
3.1
30
The OpenCL Specication . . . . . . . . . . . . . . . . . . . . . .
2.2.5
2.4
28
OpenCL - Open Computing Language . . . . . . . . . . . . . . .
2.2.4
2.3
24
26
50
52
Introduzione alla progettazione:
framework di OpenCL e librerie IVLLIB . . . . . . . . . . . . . .
53
3.1.1
Le librerie IVLLIB . . . . . . . . . . . . . . . . . . . . . .
53
Requisiti e speciche del framework . . . . . . . . . . . . . . . . .
54
3.2.1
Requisiti . . . . . . . . . . . . . . . . . . . . . . . . . . . .
54
3.2.2
Speciche . . . . . . . . . . . . . . . . . . . . . . . . . . .
55
6
INDICE
3.3
3.4
3.2.2.1
Astrazione delle meccaniche di OpenCL . . . . .
3.2.2.2
Oggetti . . . . . . . . . . . . . . . . . . . . . . .
56
3.2.2.3
Singleton Design Pattern
56
3.2.2.4
Kernel OpenCL
. . . . . . . . . . . . . . . . . .
57
3.2.2.5
Politiche di gestione . . . . . . . . . . . . . . . .
57
3.2.2.6
Gestione degli errori . . . . . . . . . . . . . . . .
58
3.2.2.7
Trasparenza diversa per utenti diversi . . . . . .
58
Diagrammi di usso e UML . . . . . . . . . . . . . . . . . . . . .
59
3.3.1
Flow Chart
60
3.3.2
Diagramma UML delle classi
3.3.3
Note progettuali
. . . . . . . . . . . . . . . . . . . . . . . . . .
. . . . . . . . . . . . . . . .
60
. . . . . . . . . . . . . . . . . . . . . . .
62
Design di Algoritmi di Image Processing con il framework di
OpenCL . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
63
3.4.1
63
3.4.2
Programmatore della libreria: Gamma Correction
3.4.3
. . . .
3.4.1.1
La Gamma Correction . . . . . . . . . . . . . . .
64
3.4.1.2
Progettazione Kernel
. . . . . . . . . . . . . . .
64
3.4.1.3
Possibile Codice del Programma . . . . . . . . .
65
Utilizzatore del Framework di OpenCL:
Filtro di Smoothing
3.5
. . . . . . . . . . . . .
. . . . . . . . . . . . . . . . . . . . .
3.4.2.1
Filtro di Smoothing
3.4.2.2
Progettazione Kernel
3.4.2.3
67
. . . . . . . . . . . . . . . .
67
. . . . . . . . . . . . . . .
68
Possibile Codice del Programma . . . . . . . . .
69
Conclusioni . . . . . . . . . . . . . . . . . . . . . . . . . .
71
Possibilità di modica ed evoluzione del progetto . . . . . . . . .
71
4 Conclusioni
4.1
56
72
OpenCL: perchè
. . . . . . . . . . . . . . . . . . . . . . . . . . .
72
4.1.1
Ecienza Computazionale . . . . . . . . . . . . . . . . . .
72
4.1.2
OpenCL è Open
73
4.1.3
Con questo framework, OpenCL è facile . . . . . . . . . .
. . . . . . . . . . . . . . . . . . . . . . .
73
4.2
Cosa porta di nuovo
. . . . . . . . . . . . . . . . . . . . . . . . .
73
4.3
Note nali . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
74
0.2 Introduzione
Questa tesi ha lo scopo di analizzare tutte le sfaccettature di un argomento che
negli ultimi anni ha visto crescere la sua rilevanza in ogni ambito applicativo:
il calcolo GPGPU.
2
Negli ultimi anni il calcolo tradizionale si è rivelato insuciente a soddisfare la richiesta di potenza necessaria all'elaborazione, in real-time o meno, di
algoritmi sempre più sosticati e mole di dati di qualità e complessità sempre
maggiori. Nonostante l'evoluzione ingegneristica produca ogni anno processori
sempre più potenti, da subito è stato percepito il bisogno dello sviluppo di un
2 GPGPU:
General-purpose computing on graphics processing units.
7
INDICE
componente a se stante dedicato alla realizzazione di calcoli specializzati su dati
di uno specico tipo.
Nei computer di tutti giorni è cresciuto un componente di importanza sem-
3 nel 1981
pre maggiore: la scheda video. Fin da quando è nata, (dalla MDA
4
allo standard SVGA [1]) la scheda video è stata costruita per permettere l'esecuzione di calcoli in parallelo necessari alla realizzazione di complessi programmi
graci. Negli anni questo componente è stato aggiornato e potenziato, versione
dopo versione, arrivando no ad oggi dove con poche decine di euro è possibile
comprare un processore graco con ottime prestazioni.
Da questa presa di coscienza nasce la consapevolezza di poter sfruttare
questo componente non solo per lo scopo pressato, ovvero come motore della
5
pipeline di rendering , ma anche come piattaforma di calcolo generico.
Il 15 febbraio 2007, NVIDIA rilascia il primo SDK per CUDA[3].
Quasi
un anno dopo nel Dicembre 2007 AMD, dopo l'acquisizione di ATI avvenuta
6
circa un anno prima , rilascia lo Stream Computing SDK[2]: l'era del calcolo
GPU-based è incominciata.
La Apple denisce un gruppo di lavoro con AMD, IBM, Intel e NVIDIA ed
inizia lo sviluppo di OpenCL[5]. Bisogna aspettare il 16 Luglio 2008 per veder
7
il gruppo della Khronos [5], composto dai maggiori esponenti nel campo fra cui
costruttori di CPU, GPU e processori embedded, unirsi per creare una specica
di un linguaggio per il calcolo parallelo.
Il 18 Novembre 2008 viene rilasciata la specica 1.0 delle OpenCL.
Obiettivo nale di questa tesi è la progettazione e la programmazione di
un framework che possa astrarre molte delle funzioni della specica OpenCL
e che sia in grado di integrarsi con delle librerie di image processing IVLLIB
8
fornendo al programmatore uno strumento facile quanto potente per realizzare
i suoi algoritmi sull'architettura OpenCL.
0.2.1 Struttura del documento
La prima parte di questa tesi analizzerà l'architettura logica delle GPU, l'esame
di una pipeline di rendering e inne verrà preso come caso di studio un esempio
di scheda video recente (NVIDIA GTX 285[6, 7]). Verranno inoltre analizzate
le principale analogie e dierenze architetturali fra GPU e CPU.
Nella seconda parte verrà presentato il linguaggio OpenCL partendo da una
descrizione generale no ad arrivare all'analisi di alcune funzioni della specica
1.0 di OpenCL. Verranno inoltre analizzate le sue caratteristiche e le logiche in
cui opera e verrà mostrato un esempio di un programma completo scritto in C
e OpenCL.
3 MDA: Monochrome Display Adapter.
4 SVGA: Super Video Graphics Array.
5 Verrà spiegato in dettaglio nelle sezioni successive.
6 Acquisizione di ATI da parte di AMD: 24 Luglio 2006[4].
7 Khronos Compute Working Group.
8 L'integrazione e il commento di queste librerie compongono
il terzo capitolo di questa tesi.
INDICE
8
Nella terza parte di questa tesi si discuteranno la progettazione e le scelte
implementative del framework realizzato in C++ e OpenCL e verranno analizzati alcuni esempi di possibili utilizzi per la programmazione di algoritmi di
image processing in OpenCL da parte di diversi utenti.
Inoltre verrà descritta la libreria IVLLIB per comprendere l'integrazione di
questo framework con la libreria.
Inne l'ultimo capitolo accoglierà le conclusioni a cui lo studio di OpenCL e
la progettazione del framework sono arrivate.
Capitolo 1
Architettura delle GPU
In questo capitolo viene analizzata l'architettura logica delle schede video, modellizzando gli esempi, per spiegare come è costruito questo componente presente
ormai in ogni computer in vendita oggi. Verranno analizzate le caratteristiche
comuni delle architetture GPU per poi entrare nel caso specico di una scheda
video recente, la GeForce GTX 285 [6, 7].
1.1 La storia delle GPU: radici ed evoluzione
Dalla realizzazione dei primi computer con terminale a schermo si è sentita
la necessità di progettare un componente staccato ed indipendente dal processore che fosse responsabile della visualizzazione sul display delle operazioni che
avvenivano nella macchina.
Una GPU
1 è un processore collegato alla scheda video creato per compiere
questo lavoro in maniera autonoma risperro alla CPU. Un accelleratore graco
contiene microchip costruiti appositamente per la realizzazione delle comuni
operazioni matematiche in oating point usate nel rendering graco. [8]
1.1.1 Evoluzione delle GPU
2
Il primo esempio di GPU si può trovare nel chip realizzato dalla ANTIC e dalla
CTIA negli anni 70: il chip svolgeva il compito di controllore hardware per la
modalità graca e testuale, la posizione delle sprite
3 e il display delle immagini a
schermo. Questo chip era il primo esempio di processore dedicato esclusivamente
al mapping del testo e della graca a schermo.
L' IBM Professional Graphics Controller fu il primo vero esempio di processore 2D/3D disponibile per i pc IBM: nel 1984, costava circa 4500$ e questo,
1 GPU: Graphics Processing Unit, unità di calcolo graco.
2 Viene mantenuta la stessa formattazione della fonte [8] da cui si ispira tutta questa sezione.
3 In graca informatica, gura bidimensionale che può essere spostata rispetto allo sfondo.
9
CAPITOLO 1.
10
ARCHITETTURA DELLE GPU
complice la mancata compatibilità con i sistemi del tempo, non permisero la
diusione di questa scheda.
L'Amiga Commodore fu il primo computer di massa a montare una scheda
dedicata esclusivamente al calcolo graco ed alla implementazione di base delle
primitive 2D via hardware.
Nel 1991, la S3 Graphics fu la prima a costruire un chip ad alte prestazioni
per la graca 2D: da allora ebbe inizio, in tutto il corso degli anni 90, la realizzazione di schede video dalle caratteristiche di volta in volta migliori: una
gara di prestazioni tutt'oggi ancora aperta.
prime API
In quegli anni si svilupparono le
4 grache ed i primi sistemi operativi con una GUI5 . A metà degli
anni novanta le schede video dovevano essere in grado di compiere calcoli per
la graca 3D sia per il computer che per le console. Sempre all'inizio degli anni
novanta nascono le OpenGL e qualche anno dopo le DirectX: le schede video
devono essere in grado di soddisfare i requisiti e le richieste di un mercato in
continua evoluzione.
Dagli anni 2000 ci fu l'avvento più importante nel campo della computer
graphics: la programmazione shader. Gli shader sono un insieme di istruzioni
software che permettono di calcorare eetti di rendering su hardware graco con
6
grande essibilità[19] . Le schede video, dopo la crescita di OpenGL e DirectX
come librerie grache, devono quindi adattarsi alle richieste del mercato e supportare shading programmabili per permettere ai programmatori di accedere alle
funzionalità che abilitano il render della scena sullo schermo. NVIDIA fu la pri-
7
ma a costruire una scheda video con shader programmabili . Pochi mesi dopo,
con l'introduzione della ATI Radeon 9700, viene creato il primo accelleratore
delle librerie Direct3D 9.0 in grado di compiere loop e lunghi calcoli matematici
in oating point grazie a pixel e vertex shader: queste GPU sono estremamente
più veloci del processore nella gestione di operazioni su immagini immagazzinate
in memoria come array.
1.1.2 La situazione odierna
Oggi, le GPU sono un componente fondamentale in ogni computer. Lo sviluppo
della graca real time o batch rappresenta una fetta del mercato dell'informatica rilevante.
Le due aziende leader nel settore sono NVIDIA e ATI/AMD.
Queste due case, come scritto prima, hanno una lunga storia nella produzione
di chip graci e ogni anno producono set di schede video sempre più potenti e
8 e delle speciche
più economiche. Lo sviluppo delle nuove librerie DirectX 11
9
OpenGL 4.0 richiedono, al giorno d'oggi, che le schede video siano in grado
di supportare calcoli sempre più complessi.
Con poche centinaia di dollari è
4 API: Application Programming Interface.
5 GUI: Graphics user interface.
6 Esistono più tipi di shader e sarebbe interassante
Viene evitato in quanto non oggetto di questa tesi.
un approfondimento sull'argomento.
7 La serie GeForce 3, nome in codice NV20.
8 http://www.microsoft.com/games/en-US/aboutGFW/pages/directx.aspx
9 http://www.khronos.org/opengl/
CAPITOLO 1.
11
ARCHITETTURA DELLE GPU
Figura 1.1: Scheda Video Ati Radeon 5850[9].
possibile entrare in possesso di una scheda che no a qualche anno fa era architetturalmente impossibile da costruire. Vengono presi due esempi di scheda
video, un esemplare di ATI/AMD e un esemplare di NVIDIA.
Entrambe queste schede sono presenti sul mercato internazionale ad un prezzo di circa 250¿. Entrambe queste schede video sono compatibili con le tecnologie sopra citate e sono in grado di supportare senza problemi ogni applicazione
videoludica e non presente sul mercato in data odierna.
Inoltre, oggi, l'applicazione di schede per calcolo general purpose è sempre
più diuso in ogni ambito dove venga richiesta la possibilità di eseguire calcoli
paralleli il più velocemente possibile
10 .
1.2 Architettura generica di una GPU
In questa sezione verrà analizzata schematicamente l'architettura delle GPU e
la realizzazione di una generica pipeline di rendering.
Verranno mostrare le
analogie fra la pipeline di rendering e le caratteristiche del calcolo GPGPU.
10 Gli
ambiti in cui vengono applicati il calcolo GPGPU sono veramente vari e spaziano
dalle applicazioni grache in real time alla medicina, dalle ricerche metereologiche agli studi
dei fenomeni di evoluzione di massa.
CAPITOLO 1.
12
ARCHITETTURA DELLE GPU
Figura 1.2: Scheda Video GeForce GTX 285[10].
1.2.1 Architettura
Esistono principalmente due tipi di schede video: le soluzioni integrate nella
motherboard del computer o le schede video dedicate (come gure 1.1 nella
pagina precedente e 1.2).
Verrà approfondita l'architettura di schede video
dedicate (si ricorda che in questa sezione si fa riferimento ad un modello teorico
di scheda video, senza fare nessun riferimento a nessuna scheda in particolare).
Una scheda video dedicata è generalmente composta da
ˆ
GPU;
ˆ
Video BIOS;
ˆ
Video Memory;
ˆ
RAMDAC;
ˆ
Outputs;
ˆ
Motherboard Interface;
ˆ
Cooling Device;
ˆ
Power Demand.
11 Interamente
ispirato a [1].
11 :
CAPITOLO 1.
13
ARCHITETTURA DELLE GPU
1.2.1.1 GPU: Graphics Processing Unit
La GPU è il componente principale di una scheda video: si occupa delle operazioni e dei calcoli in oating point, fondamentali per il calcolo 3D. Il componente principale, e sicuramente anche il più complesso, il processore graco
si caratterizza per velocità di clock espressa in Hertz (Hz ) e per il numero di
pipelines che deniscono la potenza necessaria con cui questo componente può
eettuare i calcoli in parallelo.
1.2.1.2 Video Bios
Il Bios Video è il programma di base che governa le operazioni della scheda video
e permette l'interazione fra la macchina e la scheda gestendo tutti i particolari
dovuti alla comunicazione fra questi due soggetti. Contiene tutte le informazioni
di basso livello come le frequenze di GPU, memorie, timings e voltaggio di ogni
componente che opera nella scheda video.
1.2.1.3 Video Memory
La memoria video svolge il componente di memoria centrale per la scheda video.
Dato che la memoria deve permettere accesso al processore ed agli altri componenti, deve essere il più veloce possibile: per questo vengono usate speciali
memorie high-speed o multi-port
12 . Dal 2003 a oggi, sono state costruite schede
13 sempre più veloci.
video con memorie DDR
La tabella qui sotto riassume la velocità di accesso e di trasmissione delle
memorie usate nelle schede video[1]:
Type
Memory clock rate (MHz)
Bandwidth (GB/s)
1.2 - 30.4
DDR
166 - 950
DDR2
533 - 1000
8.5 - 16
GDDR3
700 - 2400
5.6 - 156.6
GDDR4
2000 - 3600
128 - 200
GDDR5
3400 - 5600
130 - 230
1.2.1.4 RAMDAC14
La RAMDAC è un componente analogico/digitale che ha il compito di regolare
il funzionamento della scheda video rispetto allo schermo. Converte il signale da
digitale ad analogico per permettere ai dispositivi video come schermi CRT una
corretta visualizzazione dell'immagine. In seguito alla diusione di schermi per
computer digitali e alla integrazione di questo componente nel die della GPU, la
RAMDAC è andata sparendo come componente a se stante nelle schede video.
12 VRAM, WRAM, SGRAM, etc.
13 Non viene specicata la denizione
tesi; per informazioni dettagliate:
14 Random
di memorie DDR, in quanto non oggetto di questa
http://en.wikipedia.org/wiki/Double_data_rate
Access Memory Digital-to-Analog Converter.
CAPITOLO 1.
14
ARCHITETTURA DELLE GPU
1.2.1.5 Output
Nelle schede video odierne è possibile trovare uno o più fra i seguenti connettori
15 :
per collegare uno o più schermi alla scheda video
ˆ
VGA;
ˆ
DVI;
ˆ
VIVO (Video In Video Out);
ˆ
HDMI.
1.2.1.6 Motherboard Interface
L'interfaccia di comunicazione fra scheda video e scheda madre è un componente
fondamentale, sia ai ni del l'uso GPGPU
16 , sia per gli scopi standard della
GPU. Un bus veloce permette una comunicazione agile fra sistema e scheda;
oggi il bus più usato è il PCI Express 16x 2.0
17 in continua evoluzione.
1.2.1.7 Cooling device
Come ogni componente del computer, anche la scheda video (ed in particolare
la GPU) ha bisogno di dissipare il calore che produce in seguito all'uso.
Per
18 o di
19
20
un meccanismo
che mantenga la temperatura del chip sotto un certo valore .
questo motivo ogni modello di scheda video è provvisto di un dissipatore
1.2.1.8 Power Demand
Dato l'aumento di prestazioni delle schede video, col passare del tempo e della
continua evoluzione dei modelli è aumentata anche la richiesta in termini energetici. Pochi modelli riescono a essere alimentati dalla corrente presente sul bus
di comunicazione della scheda madre ed ormai ogni modello presente oggi sul
mercato dispone dei connettori da collegare direttamente all'alimentatore per
fornire la corrente necessaria ad ogni componente della scheda.
1.2.1.9 Osservazioni
Dalla descrizione precedente, è possibile intuire la presenza di una relazione
fra l'architettura di un intero computer (scheda madre, processore e memoria
centrale) e l'achitettura di una scheda video. Questa relazione è indice di due
15 Non
è interesse in questa tesi entrare nello specico dei vari connettori e delle loro
dierenze.
16 Questo dettaglio
17 PCIe x16 2.0:
sarà chiarito più avanti.
Width
(bits)=1
Ö
16,
Bandwidth(MB/s)=8000 / 16000, Type=Serial.
Clock
Rate
(MHz)=5000
18 Denito attivo se ci sono ventole o passivo senza ventole.
19 Dissipazione ad aria; sistema a liquido; dissipatore con ventole.
20 Circa 100 gradi centigradi, dipendente dal chip e dalla sua qualità.
/
10000,
CAPITOLO 1.
15
ARCHITETTURA DELLE GPU
principali fattori: la complessità necessaria a sviluppare un hardware dedicato sempre più potente e indipendente dalla macchina sottostante e la qualità
ingegneristica raggiunta nei nostri giorni.
1.2.2 Pipeline di rendering[11, 12]
Nella computer graphics, la pipeline di rendering è quel processo svolto dalla
scheda video in grado di renderizzare a schermo, quindi in un ambiente 2D, la
rappresentazione di oggetti descritti matematicamente in una scena tridimensionale; vengono quindi calcolate luci, ombre, la posizione dell'osservatore ed il
suo orientamento.
Verrà quindi descritto un approcio ad alto livello in grado di evidenziare i
passi fondamentali di cui questo processo è composto per capire come le GPU
possano essere utilizzate dai programmatori come motori per calcoli paralleli.
Il processo della pipeline di rendering è generalmente suddiviso in
ˆ
Trasformation;
ˆ
Per-Vertex Lightning;
ˆ
Viewing Trasformation;
ˆ
Primitives Generation;
ˆ
Projection Trasformation;
ˆ
Clipping;
ˆ
Rasterizazion;
ˆ
Texturing, Fragment Shading;
ˆ
Display.
21 :
1.2.2.1 Trasformation
Una GPU specica un oggetto in relazione alla sua posizione nell'ambiente in
un certo spazio di coordinate; questa operazione risulta comoda alla scheda per
l'organizzazione dei dati nella rappresentazione logica di un ambiente. Prima
di eettuare il rendering occorre però eettuare una serie di trasformazioni per
costruire l'ambiente in un sistema di coordinate comuni. Queste trasformazioni
22 . Queste operazioni23 portano
sono limitate a rotazioni, traslazioni e scaling
le schede video a compiere calcoli vettoriali su matrici e il bisogno di compiere
una quantità elevata di operazioni oating point ha richiesto che la scheda video
elabori queste operazioni con grande ecienza. Citando David Luebke e Greg
Humphreys[12]:
21 Ci sono dierenza fra la pipeline di rendering in tempo reale o quella batch
22 Esistono molte trasformazioni che non vengono citate qui in quanto non oggetto
analisi.
23 Come,
ad esempio, la rappresentazione degli oggetti in coordinate omogenee.
di questa
CAPITOLO 1.
16
ARCHITETTURA DELLE GPU
The need for ecient hardware to perform oating-point vector
arithmetic for millions of vertices each second has helped drive the
GPU parallel-computing revolution.
1.2.2.2 Per-Vertex Lightning
Dopo aver mappato ogni gura nel sistema di coordinate omogenee, la scheda
24 .
video deve occuparsi di colorare gli oggetti in base alle luci presenti nel mondo
La GPU gestisce più luci e calcola il contributo di ognuna di queste in ogni
punto
25 . Per adempiere a questo compito basta sapere che, ancora una volta,
la GPU dovrà compiere un numero elevato di calcoli vettoriali.
1.2.2.3 Viewing Trasformation
In questa parte, la scheda deve compiere calcoli per trasportare l'ambiente tridimensionale nel sistema di riferimento della camera, rispetto al punto di vista
dell'osservatore. Anche qui, una grande mole di dati vengono elaborati con calcoli matriciali e vettoriali, sviluppati sempre via hardware per avere un'ecienza
più alta possibile.
1.2.2.4 Primitives Generation
In questo processo vengono semplicemente creati i poligoni in base alle regole di costruzione dei vertici e vengono rimappate le primitive generate dalle
trasformazioni eettuate.
1.2.2.5 Projection Trasformation
La trasformazione prospettica consiste nel trasformare un ambiente tridimensionale in una gura bidimensionale rispetto al punto di osservazione della camera virtuale. Anche qui la GPU ha più modalità per determinare come gli oggetti
devono essere portati dallo spazio 3D ad una rappresentazione bidimensionale
dipendente dal punto di vista da cui la scena viene osservata.
1.2.2.6 Clipping
Ora la scheda video deve calcolare le coordinate che niscono fuori dalla viewing
frustum
26 ed deve ignorare tutti i vertici che non appartengono all'insieme.
Questa operazione accellera il processo successivo di rasterizzazione eliminando
le porzioni dell'immagine non necessarie.
24 Non
viene specicato, in quanto non oggetto di questa sezione, a quali modelli di
illuminazione stiamo facendo riferimento.
25 In
questa sezioni non ci interessa nello specico la matematica usata; l'unico nostro
interesse è la modalità di calcolo, ovvero calcolo vettoriale.
26 Potrebbe
essere tradotto nestra di visione e corrisponde alla zona che la telecamera
vede realmente e quindi non interessa analizzare i vertici che niscono fuori da questa zona.
CAPITOLO 1.
ARCHITETTURA DELLE GPU
17
1.2.2.7 Scan Conversion, Rasterization
E' sicuramente possibile che ogni poligono nella gura incroci uno o più pixel:
determinare quali sono i pixel appartenenti alla gura e quali non lo sono è
compito del processo di rasterizzazione.
In questo processo, l'immagine 2D
viene direttamente convertita in un'immagine raster.
Dato che ogni pixel è
indipendente dagli altri, le architetture delle schede video sono state concepite
con l'obiettivo di compiere queste operazioni in parallelo. Questa osservazione
ha portato a costruire schede con più pipelines che lavorano in parallelo, una
indipendentemente dall'altra.
1.2.2.8 Texturing, fragment shading
In questo stadio vengono determinati i colori dei pixel mediante interpolazione
fra questi, prendendo il colore dalla matrice raster o da texture presenti in
memoria. Le GPU salvano queste texture in memorie ad alta velocità; la scheda permette quindi ad ogni pixel di accedere indipendentemente alla memoria
texture.
Dato che l'accesso a questa memoria è un comportamento modale,
sono state sviluppare tecniche di caching per permettere che questa funzione
operi con grande ecienza nascondendo così la latenza che può avere l'accesso
in memoria.
1.2.2.9 Display
La matrice di pixel colorata può essere nalmente mandata allo schermo, o al
dispositivo di visione, per essere visualizzata.
1.3 Un esame di una scheda: NVIDIA GeForce
GTX 285
27
La premessa che occorre fare riguarda la scelta di questa sezione: per spiegare
la tecnologia OpenCL è utile entrare nel dettaglio di un esempio di architettura
di scheda video. La scelta è caduta sulla GeForce GTX 285 (immagine 1.2 nella
pagina 12) per semplici ragioni:
ˆ
è una scheda abbastanza nuova e permette l'esecuzione delle tecnologie
CUDA[3] e OpenCL;
ˆ
è la scheda su cui è stato trovato più materiale, ispirandomi prevalentemente ai tutorial distribuiti sul sito
28 ;
http://www.macresearch.org/[14]
da David Gohara[13]
ˆ
grazie alla tecnologia CUDA[3], è stato più facile reperire informazioni a
riguardo dell'architettura NVIDIA.
27 Questa sezione è pesantemente ispirata a [13].
28 A cui viene dedicato un ringraziamento speciale.
CAPITOLO 1.
18
ARCHITETTURA DELLE GPU
1.3.1 Descrizione generale
Uscita nel 15 Gennaio 2009[6] al prezzo di 340$, questa scheda viene costruita con processo produttivo a 55 nanometri. Vengono riportate direttamente le
speciche dal sito del produttore[7]:
Dato
Valore
Memory Specs:
Memory Clock (MHz)
1242
Standard Memory Cong
1 GB GDDR3
Memory Interface Width
512-bit
Memory Bandwidth (GB/sec)
159.0
Feature Support:
NVIDIA SLI Ready
2 way/3 way
NVIDIA 3D Vision Ready
yes
NVIDIA Pure Video Technology
HD
NVIDIA PhysX Ready
yes
NVIDIA CUDA Technology
yes
Microsoft DirectX
10
OpenGL
2.1
Certied for Windows 7
yes
Maximum Digital Resolution
2560x1600
Maximum VGA Resolution
2048x1536
Standard Display Connectors
HDTV Two Dual Link DVI
Multi Monitor
yes
HDCP
yes
HDMI*
Via adapter
Audio Input for HDMI
SPDIF
Thermal and Power Specs:
Maximum GPU Temperature (in C)
105 C
Maximum Graphics Card Power (W)
204 W
Minimum System Power Requirement (W)
550 W
Supplementary Power Connectors
6-pin x2
1.3.2 Dettagli dell'analisi
Fra tutte le parti di cui si compone la scheda video, verrà analizzato solo il
processore graco per capirne l'architettura e comprendere i dettagli di cui una
GPU è composta. Viene svolto quindi un approfondimento sulla struttura del
processore e viene fatto riferimento alla memoria all'interno della GPU. I disegni e le ragurazioni (a meno delle immagini 1.3 nella pagina successiva, 1.4
nella pagina 20 e 1.5 nella pagina 21) sono astrazioni molto lontane dalla realtà
implementativa del chip. Inoltre viene mostrata una dierenza qualitativa
una CPU ed una GPU.
29 Verranno
mostrate una foto di una CPU e una foto di una GPU.
29 fra
CAPITOLO 1.
19
ARCHITETTURA DELLE GPU
Figura 1.3: Processore Intel Core 2 Duo[18].
1.3.3 Dierenze qualitative fra GPU e CPU
Come si può vedere dalle immagini 1.3 e 1.4 nella pagina successiva, la CPU
(immagine 1.3) è composta principalmente di due parti: la parte dedicata all'elaborazione (in gura nella parte superiore) contenente tutti i meccanismi
30 ) e la memoria
necessari allo svolgimento delle operazioni basilari (ALU, ...
cache che occupa più di metà del chip (nella parte inferiore dell'immagine).
La memoria cache serve principalmente per mascherare le latenze delle operazioni da eettuare con la memoria centrale di un computer (accesso in lettura
e scrittura dalla ed alla memoria).
Se osserviamo piu attentamente l'immagine di una GPU (immagine 1.4 nella
pagina successiva) possiamo vedere come manchi quasi totalmente la memoria
di tipo cache ed invece il die sia completamente coperto da processori ai lati (ai
quattro angoli), e di unità di calcolo in mezzo al die.
Viene mostrata un'altra gura che permette di capire come è suddiviso il
chip della GeForce GT200 (gura 1.5 nella pagina 21).
Si nota subito la quantità impressionante di processori agli angoli della scheda video: verrà analizzato in seguito come questi processori sono logicamente
30 Non
vengono esplicitati tutti i componenti in quanto questa tesi, ed in particolare questa
sezione, non hanno l'intenzione di spiegare specicatamente le dierenze architetturali fra
CPU e GPU, ma piuttosto quella di dare un'idea sulle dierenze logiche e qualitative fra i due
oggetti.
CAPITOLO 1.
20
ARCHITETTURA DELLE GPU
Figura 1.4: Die di una GeForce GT200[15].
organizzati. Grande parte del chip è dedicata al Frame Buer (ai lati, in azzurro), componente responsabile di rendere a schermo il contenuto del buer di
memoria che contiene un frame completo di dati ed immagine. In giallo sono
evidenziate le unità ROP, i Raster Processor Pipeline. Nella pipeline di rendering, la pixel pipeline prende le informazioni dei pixel e dei texel e le processa
con operazioni di matrici e vettori no al valore del pixel.
I ROP compiono
questa transazione dai buer nella memoria locale e questo include la scrittura
e la lettura dei valori, insieme alle tecniche per miscelare questi dati[21]. Inne
in viola sono evidenziare le Texture Processing Unit che hanno il compito di
registrare i dati ed eseguire le operazioni riferite al texturing degli oggetti.
Al centro troviamo uno dei componenti più importanti, e più caratteristici, di
una GPU che consiste nel Thread Manager: le schede video, e questa è una dif-
31 responsabile
ferenza abissale rispetto ai processori, hanno hardware dedicato
della gestione dei thread.
Da questo possiamo capire che le GPU svolgono operazioni legate al multithreading come context switching e gestione delle priorità fra thread completamente via hardware e quindi con una ecienza veramente alta.
Poco sopra, in arancione, troviamo il setup raster, il processore responsabile
dell'operazione di rasterization, l'ultima parte della pipeline di rendering.
31 Nell'immagine
1.5 nella pagina successiva corrisponde all'area colorata in grigio.
CAPITOLO 1.
Figura 1.5:
ARCHITETTURA DELLE GPU
21
Die di una GeForce GT200[16] che evidenzia ogni particolare
architetturale.
CAPITOLO 1.
22
ARCHITETTURA DELLE GPU
1.3.3.1 Conclusioni dell'analisi fra CPU e GPU
Questa disserzione fra CPU e GPU serve solo a sottolineare la diversa natura
architetturale fra i due chip, talmente diversa da giusticare i risultati che si
ottengono su una GPU rispetto ad una CPU durante certe elaborazioni.
In proposito occorre citare il lavoro del collega Pigazzini Andrea[30] che ha
eettuato nell'anno 2007/2008 la tesi dal titolo Utilizzo di CUDA nell'imple-
mentazione di algoritmi di supporto all'Elaborazione delle Immagini . Il lavoro
del collega era completamente incentrato sullo studio delle prestazioni di alcuni
algoritmi di image processing eseguiti sulla scheda video rispetto che sul processore centrale e i suoi risultati sono stati sconcertanti: con una GPU da 100¿
le prestazioni erano totalmente a favore della GPU in ogni test.
1.3.4 Il chip GT200: organizzazione logica
Quì di seguito verrà mostrata un modello teorico necessario per capire i dettagli
fondamentali, comuni comunque a tutte le GPU, del chip NVIDIA GTX285
derivante dal chip GT200[6].
Negli schemi sono stati riportati solo i particolari interessanti all'analisi del
chip, particolari che avranno ancora più senso quando verrà spiegata la logica
di OpenCL.
L'immagine 1.6 nella pagina seguente rappresenta un modello della GPU
della GTX285
32 . Il miglior modo per descriverla è citare la fonte [22]:
There are no two ways about it Nvidia's GT200 GPU is an
absolute brute. It's manufactured using TSMC's 65nm process, features approximately 1.4 billion transistors and packs a total of 240
thread processors running at 1,296MHz (on the GeForce GTX 280)
the result is a GPU that delivers 933.12 gigaFLOPS of compute
power at peak. These numbers alone make it the largest and most
complex GPU ever to be made.
Dalla gura 1.6 nella pagina successiva è possibile osservare subito come è organizzata la struttura generale della GPU. Verranno analizzate solo le caratterisriche di interesse per poter poi comprendere l'architettura logica su cui opera
OpenCL.
Il chip GT200 è composto
33 da 10 TPC, Thread Processing Cluster34 . Ques-
ta è la prima divisione gerarchica che abbiamo all'interno della GPU e d'ora in
poi parlemo di come un TPC è suddiviso al suo interno.
In particolare un
Thread Processing Cluster è diviso al suo interno in tre SM, gli Streaming
Multiprocessor.
32 Il
modello reference della citazione [23] fa riferimento alla GTX280 che a livello architet-
turale è praticamente identica alla GTX285; inoltre essendo un modello teorico questi due
chip sono concettualmente uguali.
33 Nella gura 1.6 nella pagina seguente i rettangoli verdi grandi.
34 Questa serie di schede video è stata chiamata la serie 10 probabilmente
da 10 TPC rispetta ai soli 8 della serie precedente, la serie g92[23].
perchè costituita
CAPITOLO 1.
ARCHITETTURA DELLE GPU
23
Figura 1.6: Architettura logica della GPU GTX285[23].
Figura 1.7:
GT200[23].
Architettura logica di un TPC della serie g92 e uno della serie
CAPITOLO 1.
ARCHITETTURA DELLE GPU
24
Figura 1.8: Particolare di un TPC con memoria cache[23].
Uno Streaming Multiprocessor ha una memoria texture L1 cache condivisa
di 24 KByte.
Questo genere di memoria permette un accesso read/write molto veloce e
condiviso fra gli Streaming Multiprocessor.
In gura 1.9 nella pagina succes-
siva, il SMC è lo Streaming Multiprocessor Controller, il responsabile della
suddivisione del lavoro fra più Streaming Multiprocessor.
1.3.4.1 Streaming Multiprocessor
Scendendo nel dettaglio, evitando i particolari non utili all'analisi in corso, uno
Streaming Multiprocessor (in gura 1.9 nella pagina seguente) è a sua volta
suddiviso in:
35 ;
ˆ
8 Streaming Processor
ˆ
2 Special Function Unit;
ˆ
1 Double Precision Unit;
ˆ
16 KB di memoria locale (shared memory) condivisa fra gli otto Streaming
Processor;
Gli 8 Streaming Processor sono le unità di base, quelle che eseguono il calcolo.
Vengono ben schematizzate dalla gura 1.10 nella pagina 26. Per una chiara e
breve descrizione cito la fonte [23]:
35 Chiamati anche streaming cores,
NVIDIA) o, generalmente, cores.
CUDA core (dato che stiamo analizzando un'architettura
CAPITOLO 1.
ARCHITETTURA DELLE GPU
Figura 1.9: Particolare di uno Streaming Multiprocessor[23].
25
CAPITOLO 1.
ARCHITETTURA DELLE GPU
26
Figura 1.10: Schema logico di uno Streaming Processor.
In sintesi uno Stream Processor o Processor Core è un microprocessore completo dotato di due ALU e una FPU, non ha cache e
la sua dote è quella di ripetere tonnellate di operazioni matematiche!
Un SP spende gran parte del suo tempo lavorando sui pixel o vertex
data e non è importante il fatto che non abbia cache. Un solo SP è
abbastanza inutile, ma se inserito insieme ad altri in unità come le
SM (Streaming Multiprocessor) si inizia a sentire la sua potenza.
Le due Special Function Unit sono delle unità responsabili di compiere operazioni e calcoli particolari: funzioni trigonometriche, esponenziali e logaritmiche.
Il fatto che venga dedicato dell'hardware specico permette a queste
operazioni, frequenti nel calcolo della pipeline di rendering, di essere svolte con
grande ecienza.
L'unità a doppia precisione ha il compito di svolgere calcolo in virgola mobile
e di gestire dati molto più complessi e precisi di quelli in oating point con
precisione singola, necessario in alcuni tratti della pipeline ed utile nel calcolo
general purpose.
La memoria locale di 16 KB è la memoria più veloce che uno Streaming Processor può usare. Sfruttata in un certo modo, questa memoria ha prestazioni elevatissime e permette agli Streaming Processor di usarla in maniera molto simile
e con prestazioni molto vicine ai registri dei core. Esistono tecniche di programmazione in CUDA e in OpenCL che permettono di sfruttare questa memoria il
più ecientemente possibile. Un programmatore che riesce a usufruire di questo
strumento, riesce a creare del codice con performance assolutamente incredibili.
1.4 Conclusioni sull'architettura delle GPU
Ci sono 10 TPC, suddivisi in 3 SM con all'interno 8 SP. In un chip GT200
esistono quindi:
10 T P C ∗ 3 SM ∗ 8 SP = 240 Streaming P rocessor
CAPITOLO 1.
ARCHITETTURA DELLE GPU
27
Un numero così elevato di core ed un'unità hardware interamente dedicata
36 vogliono suggerire una sola cosa: le schede video sono
alla gestione dei thread
state create per macinare calcoli in parallelo, in quantità normalmente non
concepibili con task multithreading svolti su CPU ordinarie. Le GPU sono state
create per gestire milioni di operazioni in parallelo: amano il context switching,
hanno dell'hardware dedicato per compiere queste operazioni e occorre creare dei
linguaggi che permettono di usare e piegare così tanta potenza all'esecuzione
di calcoli complessi, che ormai permeano ogni ambito scientico e non.
Da queste esigenze è nato il calcolo GPGPU. Un nuovo modo di concepire i
problemi quotidiani ed un nuovo modo di risolverli.
36 Si
ricorda il Thread Manager in gura 1.5 nella pagina 21.
Capitolo 2
OpenCL: Open Computing
Language
In questo capitolo ci concentreremo sullo studio, sotto ogni punto di vista, del
linguaggio OpenCL[5].
Questo capitolo aronterà, usando la stessa struttura
della specica OpenCL rilasciata dalla Khronos[24], gli argomenti necessari a
spiegare cosa è, perchè viene usato, come è stato pensato e come funziona il
linguaggio OpenCL. Questo capitolo trae ispirazione dalla specica del linguaggio, fonte [24], ed dalla fonte [5]. Alcuni cenni derivano dai tutorial di David
1
Gohara[13] , presi da
http://www.macresearch.org/opencl[14].
Questo capitolo è suddiviso in tre parti principali: un'introduzione ad OpenCL, come e perchè è stato creato, la sua storia e le motivazioni per cui è stato scelto per questa tesi.
La seconda parte parlerà dell'architettura logica di
OpenCL, i suoi modi di intedere i device, i dati ed i programmi. La terza ed
ultima parte mostrerà un esempio classico per far capire i meccanismi basilari
di OpenCL. L'esempio, commentato riga per riga, guida il lettore in un programma completo dalla fase di inizializzazione alla fase di chiusura e pulizia
dell'ambiente.
2.1 OpenCL - Open Computing Language
Questa sezione ha il compito di introdurre il linguaggio OpenCL e spiegare le
scelte basilari per cui è stato costruito e per cui è stato scelto per programmare
il framework in oggetto di questa tesi.
2.1.1 Storia
Il 15 Febbraio 2007 la prima beta di CUDA veniva rilasciata per i sistemi Microsoft Windows e Linux[3]; sempre nel 2007, a Dicembre, ATI/AMD rilascia il
1A
cui viene dedicato un ringraziamento speciale in questa tesi.
28
CAPITOLO 2.
SDK
OPENCL: OPEN COMPUTING LANGUAGE
29
2 di ATI Stream[2].
Apple incomincia a sviluppare OpenCL e, dopo una collaborazione con
3
le maggiori industrie nel campo , propone la realizzazione della specica alla
Khronos Group. Il 16 Giugno 2008, viene formato il Khronos Compute Working Group con la partecipazione dei maggiori rappresentanti nel campo delle
CPU, GPU, processori embedded e compagnie software.
Dopo cinque mesi di lavoro questo gruppo crea la specica 1.0 di OpenCL,
il 18 Novembre 2008. Il giorno 8 Dicembre 2008 viene rilasciata al pubblico la
specica.
2.1.2 Cos'è
OpenCL è un framework per scrivere programmi che possano essere eseguiti su
un insieme eterogeneo di piattaforme fatte di CPUs, GPUs e altri processori.
OpenCL è un linguaggio, basato sul C99, per scrivere funzioni speciali (kernel)
da eseguire sui device OpenCL e per creare un ambiente in cui queste funzioni
possano lavorare. OpenCL permette di scrivere istruzioni con un altissimo tasso
di parallelismo di dati e di programmi.
Gestito dall'azienda no-prot Khronos, OpenCL è l'analogo delle speciche
Open come OpenGL e OpenAL.
2.1.3 Perchè
OpenCL nasce dall'esigenza di creare un linguaggio totalmente indipendente
dalla piattaforma per risolvere problemi paralleli e permettere al programmatore
di usare tutti i device all'interno di un computer senza limitazioni e senza dover
creare del codice dedicato per ogni singolo device. La portabilità del codice è
uno degli obiettivi primari di OpenCL. Ogni implementazione deve rispettare i
vincoli della specica OpenCL[24] e, se soddisfa i requisiti, il codice può essere
eseguito su qualsiasi computer che abbia installato un ambiente, software e
hardware, compatibile con OpenCL.
Un'altro vantaggio, ed uno degli obiettivi principali che OpenCL si impone, è
quello di creare un unico codice indipendente non solo dalla macchina, ma anche
dai device sottostanti. Questo vuol dire che su una macchina lo stesso codice
OpenCL può essere eseguito, senza MAI essere modicato, su uno o più device
all'interno del computer ed il codice del kernel, precedentemente compilato o
compilato in tempo reale, rimane lo stesso codice scritto una sola volta dal
programmatore senza che questo debba ricompilare il programma ogni volta
che viene cambiato un device o ne viene scelto un altro in tempo reale.
Ovviamente il programma verrà eseguito con prestazioni dipendenti dal (o
4
dai) device scelti .
2 Source Development Kit
3 AMD, Intel, Ati e NVIDIA.
4 C'è un video molto interessante
che analizza questo aspetto realizzato da ATI/AMD:
http://www.youtube.com/watch?v=7PAiCinmP9Y
CAPITOLO 2.
30
OPENCL: OPEN COMPUTING LANGUAGE
Questa caratteristica rende OpenCL uno strumento versatile, potente e agile;
rimane al programmatore creare applicativi che sfruttano queste carattestiche
al meglio delle loro potenzialità.
2.1.4 Come
Prima di entrare nel dettaglio di OpenCL, della sua architettura e delle sue
caratteristiche, introduco brevemente la logica di una computazione OpenCL.
Un programma OpenCL dovrà, in ordine:
ˆ
cercare dei device OpenCL validi;
ˆ
costruire un contesto, un ambiente che racchiude l'insieme di device,
code (queue), dati e programmi OpenCL;
ˆ
creare delle code di esecuzione;
ˆ
compilare e creare dei programmi OpenCL, compilando e creando in tempo
reale uno o più programmi composti da uno o più kernel;
ˆ
creare i dati necessari alla computazione OpenCL;
ˆ
eseguire i kernel e leggere i risultati.
Nella terza parte verrà spiegato ogni passo sopra descritto e la logica che questo
comporta.
2.1.5 Perchè OpenCL per Image Processing
Esistono due valide ragioni per cui è stato scelto OpenCL come linguaggio
per programmare questo framework e utilizzarlo per creare algoritmi di image
processing:
ˆ
OpenCL è un linguaggio che permette di risolvere problemi con alto tasso
di parallelismo e la maggior parte degli algoritmi di image processing sono
5
strutturalmente compatibili con un usso di esecuzione parallelo ;
ˆ
OpenCL, come dice la parola stessa, è Open e completamente platform
indipendent. Questo porta a scegliere OpenCL come linguaggio per realizzare un framework di elaborazione delle immagini in quanto verrà scritto
codice in grado di girare su ogni macchina che abbia un ambiente OpenCL
funzionante; infatti, al giorno d'oggi, sono stati creati ambienti pratica-
6
mente per ogni sistema operativo . Inoltre, la losoa Open di OpenCL
5 Come, perchè e quali algoritmi di image processing fanno parte di questa famiglia non verrà
arontato qui in quanto non oggetto di questo capitolo. Per ogni riferimento sull'argomento
si consiglia il lavoro del collega [30]
6 Insieme
alla tesi vengono consegnati gli eseguibili per impostare un ambiente OpenCL
su sistema operativo Windows e Linux (Debian based e Open SUSE). Per quanto riguarda i
sistemi Mac esiste una implementazione di OpenCL già installata nel sistema operativo Snow
Leopard.
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
31
si sposa perfettamente con le librerie in cui questo framework dovrà essere
7
integrato .
2.2 The OpenCL Specication
Questa sezione si ispira fortemente al documento The OpenCL Specication
1.0 [24]. Obiettivo di questa parte consiste nel dare un'idea generale dell'architettura logica con cui è stato costruito questo linguaggio. Quindi, dall'inizio
alla ne, verranno ripercorsi gli stessi passi in cui è divisa la specica di OpenCL, citando anche la fonte [25] e integrando con [13]. Dalla specica viene presa
solo la prima parte.
Verranno discussi solo gli argomenti necessari a capire un programma OpenCL in via generica, dando un ampio taglio alle informazioni non necessarie e
sottolineando le conoscenze necessarie a capire l'esempio dopo mostrato e le
scelte implementative del framework realizzato.
2.2.1 The OpenCL Architecture
OpenCL è un framework per la programmazione parallela che include un linguaggio, API, librerie ed un sistema in tempo reale per supportare lo sviluppo
del software. Per descrivere l'idea principale di OpenCL, useremo una gerarchia
di modelli:
ˆ
modello di piattaforma;
ˆ
modello di esecuzione;
ˆ
modello di memoria;
ˆ
modello di programmazione.
2.2.1.1 Modello di piattaforma
La piattaforma di OpenCL (come in gura 2.1 nella pagina successiva) consiste
in un host connesso ad uno o più OpenCL device (compute device). Un device
OpenCL è diviso in uno o più Compute Units, unità di computazione, che
sono a loro volta divise in più Processing Elements, elementi di processo. La
computazione vera e propria avviene all'interno dei Processing Elements.
Le applicazioni OpenCL mandano comandi dall'host ai Processing Elements
all'interno dei device.
Ogni Processing Element esegue un singolo usso di
instruzioni come unità SIMD
8 o come unità SPMD9 .
Se si vuole fare un'analogia fra il modello di piattaforma e un computer,
l'host può essere un programma che esegue su un computer e i compute device
7 La
losoa della libreria IVLLIB e come questo framework verrà integrato, sono oggetto
del capitolo 3.
8 Single
9 Single
Instruction, Multiple Data[26].
Process, Multiple Data[27].
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
Figura 2.1: Architettura di OpenCL
32
CAPITOLO 2.
33
OPENCL: OPEN COMPUTING LANGUAGE
possono essere il processore e/o la scheda video (ad esempio, per ricollegarci
al capitolo 1, una GTX285).
Nella GeForce GTX285, gli Streaming Proces-
sor possono essere visti come le compute units di OpenCL ed all'interno degli
streaming processor si possono identicare i thread che sono l'ultima sottodivisione dei core di una scheda NVIDIA; i thread sono i processing elements di
OpenCL.
In un Intel Core2 Duo, i due core sono le due compute units e in ogni core
10 .
ci sono i processing elements
2.2.1.2 Modello di esecuzione
L'esecuzione in OpenCL avviene in due parti: i kernels che eseguono codice sui
device OpenCL e il programma host che esegue sull'host.
Per quanto riguarda i kernel, quando uno di questi è eseguito su un device
OpenCL deve essere denito uno spazio di lavoro (o spazio di indici), ovvero
una dimensione che indichi come dividere le istanze del problema.
L'istanza
di un kernel viene chiamata work-item e rappresenza un'istanza del generico
kernel in un punto specico del problema; questo work-item è identicato da
un indice globale (e in alcuni casi anche un indice locale) che denisce di che
parte del problema fa parte. Quando viene lanciata l'esecuzione di un kernel,
vengono create
n
istanze del kernel (con
n (spazio degli indici))
ognuna per
ogni indice presente nello spazio di lavoro degli indici.
Questo approcio viene spiegato meglio con un esempio: astraendo per un
attimo da un possibile ambiente OpenCL, supponiamo di avere un immagine e
di voler fare un'operazione per ogni pixel e supponiamo che questa operazione sia
indipendente dal valore degli altri pixel. Rappresentiamo l'immagine a livelli di
grigio e quindi come una matrice in cui ogni pixel rappresenta il livello di grigio in
quella posizione. Se l'immagine è, ad esempio,
1024 ∗ 768 = 786432 pixels allora
questo problema verrà denito, in OpenCL, come un problema a due dimensioni,
in cui la prima dimensione sarà uguale a
1024 e la seconda sarà uguale a 768.
786432 work-items che eseguono
Questo vuol dire che in esecuzione ci saranno
11 ed ognuno compierà l'operazione su pixel che rappresenta.
in parallelo
Tornando all'architettura di OpenCL, i work-items sono raggruppati in workgroups.
I work-items avranno quindi un identicativo globale univoco ed un
identicativo locale al work-group univoco.
In OpenCL 1.0, lo spazio di indici supportato è chiamato NDRange.
Un
NDRange è uno spazio di indici N-Dimensionale. OpenCL 1.0 supporta spazi
di 1, 2 o 3 dimensioni.
Contesto e coda comandi
L'host denisce un contesto per l'esecuzione dei kernel. Il contesto include le
seguenti risorse:
10 Questi
riferimenti non sono precisi e non sono presi da nessuna fonte. Il loro scopo è solo
quello di rendere l'idea fra il modello di piattaforma di OpenCL e un'architettura logica di un
processore e/o di una scheda video.
11 In
relazione alla potenza e capacità del device OpenCL scelto.
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
34
ˆ
Device: la collezione di device utilizzabili dall'host;
ˆ
Kernel: le funzioni OpenCL che vengono eseguiti su dispositivi OpenCL;
ˆ
Oggetti programma: il codice e/o gli eseguibili che implementano i kernel;
ˆ
Oggetti memoria:
un insieme di oggetti memoria visibile all'host ed ai
device OpenCL; contengono i dati che saranno trasferiti ed usati sui device
OpenCL;
L'host deve inotre creare una o più code di comando (command queue) per
coordinare l'esecuzione dei kernel sui device.
Categorie di kernel
Esistono due tipi di kernel: i kernel OpenCL e i kernel nativi. I kernel OpenCL
sono i kernel scritti nel linguaggio OpenCL C e compilati con un compilatore
OpenCL. I kernel nativi sono un altro tipo di kernel, opzionali rispetto ai primi
e dipendenti dalla piattaforma
12 .
2.2.1.3 Modello di memoria
Diversamente dalla programmazione ordinaria, in OpenCL, esiste una gerarchia
di memoria che il programmatore può (e deve) controllare.
I work-items che
stanno eseguendo un kernel hanno accesso a quattro diverse zone di memoria:
ˆ
Memoria globale: questa memoria permette accessi in lettura e scrittura
a tutti i work-items in tutti i work-groups;
ˆ
Memoria costante:
questa memoria è una regione che rimane costante
durante l'esecuzione del kernel; viene inizializzata dall'host.
ˆ
Memoria locale: la memoria locale è una memoria condivisa fra i workitems in un work-group; viene usata come memoria di sincronizzazione
all'interno di un work-group;
ˆ
Memoria privata: la memoria privata è lo spazio di memoria riservato ad
ogni work-item. Ogni work-item ha la sua memoria privata e questa non
può essere condivisa con nessuno e nemmeno l'host può inizializzarla.
Per leggere e scrivere dalla e nella memoria dei device OpenCL, l'host deve
usare apposite funzioni presenti nel framework OpenCL.
Per quanto riguarda la consistenza della memoria, OpenCL non dispone di
nessun meccanismo implicito per gestire gli accessi alla e dalla memoria. Non
è quindi garantita la consistenza dello stato di una variabile condivisa fra più
work-items.
OpenCL mette a disposizione, tramite funzioni built-in, meccanismi espliciti
di sincronizzazione fra work-items in un work-group.
12 In
questa tesi non verranno trattati i kernel di questo tipo.
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
35
Figura 2.2: Architettura di memoria di un device OpenCL[28]
2.2.1.4 Modello di programmazione
Il modello di programmazione di OpenCL supporta il parallelismo di dati e il
parallelismo di processi.
La programmazione riferita al parallelismo di dati consiste in una sequenza
di istruzioni applicate ad elementi multipli di oggetti in memoria. Lo spazio di
indici denisce una esecuzione di un work-item e come i dati vengono mappati
nel work-item.
Per parallelismo di processi, OpenCL denisce un modello in cui una singola
istanza di un kernel viene eseguita indipendetemente su ogni spazio dell'indice.
Da questo modello di programmazione nasce l'esigenza di un meccanismo di
sincronizzazione. OpenCL denisce due domini di sincronizzazione:
ˆ
sincronizzazione di work-items all'interno di un work-group;
ˆ
sincronizzazione di comandi immessi in una o più code all'interno di un
contesto.
2.2.2 Il Framework OpenCL
Il framework OpenCL permette alle applicazioni di usare un host e uno o più
device OpenCL come un unico eterogeneo sistema di computer parallelo.
framework è composto da:
Il
CAPITOLO 2.
ˆ
OPENCL: OPEN COMPUTING LANGUAGE
36
OpenCL platform layer: il layer di piattaforma permette di scoprire devices e le loro caratteristiche e inizializzare contesti;
ˆ
OpenCL runtime: il runtime permette al programma dell'host di manipolare i contesti una volta creati;
ˆ
OpenCL compiler: il compilatore OpenCL crea programmi eseguibili che
contengono kernels OpenCL. Il linguaggio di programmazione OpenCL C
implementato dal compilatore supporta un sottoinsieme dell'ISO C99 con
l'estensione per il parallelismo.
2.2.3 Il livello piattaforma di OpenCL
Questa parte riguarda principalmente le funzioni necessarie ad inizializzare l'ambiente di lavoro di OpenCL. Vengono riassunti i passi fondamentali presi dalla
specica del linguaggio OpenCL[24].
Per quanto riguarda i comandi, occorre specicare che molti dei comandi di
OpenCL ritornano il controllo al programma host appena possibile e non c'è
nessuna garanzia che quando una funzione venga chiamata e ritorni il controllo
al programma, questa abbia nito la sua esecuzione. Il disegno 2.3 nella pagina
seguente chiarisce questo concetto.
Inoltre, praticamente tutte le funzioni di OpenCL ritornano un codice di
errore che dice se l'esecuzione della funzione è andato a buon ne o se c'è stato
qualche errore; in caso di errore il codice è diverso in base all'errore riscontrato
nell'esecuzione della funzione.
2.2.3.1 Richieste sulla piattaforma
Prima di tutto occorre scegliere ed inizializzare la piattaforma su cui intendiamo lavorare. In caso ce ne sia più di una, la funzione clGetPlatformIDs
ritorna una lista con tutte le piattaforme disponibili. Se volessimo conoscere più
informazioni su una piattaforma, possiamo interrogare il sistema con una query
con la funzione clGetPlatformInfo
13 .
2.2.3.2 OpenCL Device
Dopo aver scelto la piattaforma, procediamo ad una delle fasi più importanti di
un programma OpenCL: la scelta dei device. La funzione clGetDeviceIDs
ritorna la lista dei device OpenCL disponibili nel computer. OpenCL, inoltre,
permette di formulare richieste anche per device specici (solo CPU, solo GPU,
...). Anche in questo caso OpenCL mette a disposizione una funzione, clGet-
DeviceInfo, per interrogare un device sulle sue caratteristiche. Questa funzione è molto utile se si vogliono conoscere tutti i dettagli del device su cui si
sta lavorando: dalle estensioni, al supporto immagini, dal massimo numero di
work-items al massimo numero di work-group.
13 Per
eseguire in un ambiente con SDK ATI Stream è necessario scegliere una piattaforma
per la crezione del contesto OpenCL.
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
37
Figura 2.3: Esempio di usso di un comando OpenCL rispetto all'esecuzione di
una chiamata a funzione di un qualsiasi linguaggio procedurale.
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
38
2.2.3.3 Il Contesto
L'ultima parte fontamentale del livello piattaforma consiste nella creazione di
uno o più contesti.
Un constesto in OpenCL rappresenta un ambiente, un
insieme di dati, programmi (intesi come insieme di kernel) e un insieme di device
che lavorano su una piattaforma.
Con la funzione clCreateContext viene creato un contesto valido su una
piattaforma che andrà a lavorare su un insieme (una lista) di device.
OpenCL permette anche di creare un contesto specico per un solo device
con la funzione clCreateContextFromType.
Anche in questo caso OpenCL permette di avere informazioni sul contesto
con la solita procedura, svolta questa volta dalla funzione clGetContextInfo.
2.2.4 The OpenCL Runtime
Questa è la sezione principale per capire come funziona in via generale un programma OpenCL. Verranno discusse tutte le chiamate API che gestiscono le
code di comandi, gli oggetti memoria, come scrivere e leggere su questi, i kernel
e le modalità di esecuzione.
2.2.4.1 Code di comando
Gli oggetti OpenCL come oggetti memoria, programmi e kernel sono creati
all'interno di un contesto.
Le operazioni su questi oggetti sono eettuate at-
traverso l'uso di code di comando (command queue).
Le code di comando
rappresentano oggetti in cui il programmatore impila una serie di comandi e
questa tenta di risolverli il prima possibile.
Avere code multiple permette al-
l'applicazione di incodare più comandi in parallelo anche se OpenCL non ha un
meccanismo specico per la gestione della concorrenza fra code.
La funzione clCreateCommandQueue crea un oggetto coda su uno specico device. Occorre notare che su un device possono essere aperte più code ma
una coda può e deve essere aperta su uno ed un solo device OpenCL. Il solito comando clCommandQueueInfo permette di interrogare la coda sulle sue
proprietà.
2.2.4.2 Oggetti di memoria
In OpenCL esistono due tipi di oggetti di memoria:
gli oggetti buer e gli
oggetti immagine. Un buer è usato per memorizzare dati monodimensionali
mentre l'immagine è usata per creare texture, frame buer o immagini bi- o
tri- dimensionali.
Gli elementi di un buer possono essere i tipi base, i tipi
di OpenCL (come vectors, ad esempio) o struct C. Le immagini rappresentano
texture o frame-buer.
Le principali dierenze fra buer e immagine sono:
ˆ
gli elementi in un buer sono memorizzati sequenzialmente e si accede in
scrittura e/o lettura tramite puntatori. Le immagini sono immagazzinate
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
39
in un formato nascosto al programmatore che deve usare funzioni built-in
di OpenCL C per accedere in scrittura e/o lettura di un'immagine.
ˆ
Per un oggetto buer, i dati sono organizzati nello stesso modo in cui
vengono visti nel kernel mentre nel caso degli oggetti immagine il formato
con cui l'oggetto viene scritto potrebbe non essere lo stesso di quello usato
per leggere l'oggetto nel kernel.
Oggetti Buer
La funzione clCreateBuffer crea un oggetto buer di tipo cl_mem. Per
scrivere, leggere e copiare si usano i comandi clEnqueueReadBuffer, clEn-
queueWriteBuffer e clEnqueueCopyBuffer.
In questo caso, clEn-
queueReadBuffer serve per leggere un buer dalla memoria del device a quella dell'host mentre clEnqueueWriteBuffer scrive uno stream di dati dalla
memoria dell'host a quella del device. All'interno dei kernel, l'accesso ai buer
avviene tramite puntatori, con la stessa sintassi del linguaggio C. Per conoscere
le proprietà degli oggetti memoria si usa il comando clGetMemObjectInfo.
Oggetti Immagine
La funzione clCreateImage2D e clCreateImage3D creano, rispettivamente, immagini bidimensionali e tridimensionali.
Il processo di creazione di
un'immagine è più laborioso di quello del buer in quanto bisogna decidere di
quanti canali è composta l'immagine e con quanti bit rappresento ogni canale.
I dettagli non verranno elencati e si lascia al lettore la fonte [24] per maggiori
particolari.
Dato che il trattamento di una immagine, con tutte le conseguenze del caso,
non è banale, OpenCL mette a disposizione la funzione clGetSupportedIm-
ageFormat per controllare i formati di immagine supportati dal device su cui
si sta lavorando
14 .
Per scrivere, leggere e copiare immagini si usano i comandi clEnqueueRead-
Image, clEnqueueWriteImage e clEnqueueCopyImage. L'uso di queste
funzioni è analogo a quello applicato ai buer. OpenCL fornisce anche la possibilità di copiare un'immagine in un buer con la funzione clEnqueueCopyIm-
ageToBuffer. Per conoscerne le proprietà si usa il comando clGetImageInfo.
Le immagini hanno una caratteristica peculiare: la lettura (e scrittura) nel
kernel non avviene semplicemente tramite puntatore ma occorre usare funzioni
built-in di OpenCL C per leggere i dati.
Per leggere, inoltre, occorre costru-
ire un oggetto Sampler ovvero un oggetto che dice quali sono le modalità di
campionamento dell'immagine.
Questo può essere fatto nel kernel creando un oggetto Sampler costante o
nel codice host con il comando clCreateSampler da cui si possono prelevare
informazioni con clGetSamplerInfo.
14 La
CPU e la GPU, ad esempio, gestiscono un insieme di immagini di tipo diverso.
CAPITOLO 2.
40
OPENCL: OPEN COMPUTING LANGUAGE
2.2.4.3 Oggetti programma
Un programma in OpenCL è un insieme di kernel
15 e di funzioni ausiliarie
rappresentate come stringa di caratteri.
Un programma incapsula al suo interno le seguenti informazioni: un contesto,
la stringa o il binario del programma, l'ultimo build avvenuto con successo
compresi la lista dei device per cui è stato fatto il build e le opzioni di build, e
per ultimo il numero di kernel collegati al programma.
Un programma viene creato con la funzione clCreateProgramWith-
Source se abbiamo il programma sotto forma di stringa mentre, se abbiamo il
binario del programma, lo creiamo con clCreateProgramWithBinary.
16 .
Per quanto riguarda la compilazione (building) abbiamo clBuildProgram
Per avere informazioni sul programma e sulla sua compilazione (warning, errori,
...) usiamo clGetProgramInfo e clGetProgramBuildInfo.
2.2.4.4 Oggetti kernel
Un oggetto kernel è una funzione dichiarata in un programma.
Un kernel è
identicato dal qualicatore __kernel nella signature della funzione.
Per
creare un kernel da un programma dobbiamo specicarne il nome nella funzione
clCreateKernel. Se vogliamo creare tutti i kernel in un programma passiamo
un array alla funzione clCreateKernelsInProgram.
Quando creiamo un oggetto kernel, prima di usarlo dobbiamo impostare gli
argomenti che quel kernel necessita per essere eseguito. Questo compito viene
svolto dalla funzione clSetKernelArg, argomento per argomento.
Se vogliamo informazioni a proposito delle caratteristiche di un kernel possiamo interrogare il sistema con la funzione clGetKernelInfo, mentre se
vogliamo sapere i particolari di un kernel per uno specico device usiamo clGetK-
ernelWorkGroupInfo.
2.2.4.5 Esecuzione dei kernels
L'esecuzione dei kernel viene svolta dalla funzione clEnqueueNDRangeKer-
nel. Conviene spendere qualche particolare in più su questa funzione in quanto
l'esecuzione è un momento delicato in una computazione OpenCL.
Come molte funzioni OpenCL, anche questa funzione ritorna subito il controllo al usso del programma dell'host appena chiamata. Questo implica che se
abbiamo due kernel dove il secondo usa i dati del primo e nel programma questi
vengono lanciati uno di seguito all'altro, i dati saranno molto probabilmente
corrotti. Questo particolare è molto importante nell'esecuzione di code multiple
e quando abbiamo bisogno di leggere i dati subito dopo l'esecuzione del kernel.
Quando lanciamo l'esecuzione di un kernel, dobbiamo specicare, tramite
l'oggetto stesso, che kernel andiamo ad eseguire e la coda in cui impiliamo questa
operazione. Anche la scelta della coda è una scelta non univoca e non banale
15 Speciali funzioni
16 Per le opzioni di
dichiarate con l'identicatore __kernel.
building guardare la fonte [24].
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
41
in quanto apre molte scelte prograttuali in fase di costruzione di un programma
OpenCL.
Un'altra funzione particolare è clEnqueueTask che permette l'esecuzione
di un singolo kernel alla volta.
Per quanto riguarda i kernel nativi, la funzione clEnqueueNativeKer-
nel permette l'esecuzione di un kernel scritto in C/C++ non compilato con il
compilatore OpenCL.
2.2.4.6 Oggetti evento
Gli oggetti evento sono oggetti che servono ad identicare le azioni OpenCL che
si riferiscono a esecuzione di comandi o lettura, scrittura e copia su e da oggetti
di memoria. Gli eventi servono quindi a tracciare gli stati dell'esecuzione di un
comando. Quando mettiamo nella coda una esecuzione, la sua chiamata ritorna
un oggetto evento che identica quel comando.
La funzione clWaitForEvents permette di bloccare il thread host nchè
uno o più eventi non si compiono. Si può interrogare il sistema a proposito delle
informazioni di un evento con la funzione clGetEventInfo.
Dato che le code possono essere eseguite in ordine o fuori ordine, le funzioni
di OpenCL per sincronizzare gli eventi sono molto importanti. Sono quindi di
rilievo anche le funzioni clEnqueueMarker, clEnqueueWaitForEvents
e clEnqueueBarrier che metteno nella coda, rispettivamente, un marker che
ritorna un evento che può essere usato per impostare una wait su questo marker,
una lista di uno o più eventi che devono accadere prima che la coda continui
ad eseguire comandi e una barriera che blocca l'esecuzione di una coda anchè
tutti i comandi impilati prima di questo comando siano eseguiti prima che la
coda continui con i rimanenti.
Gli oggetti evento si possono anche usare per catturare le informazioni
del prolo che misura l'esecuzione del tempo di un comando con la funzione
clGetEventProfileInfo.
2.2.4.7 Flush and nish
Inne occorre analizzare due comandi molto importanti nell'esecuzione di kernels in una coda: clFlush e clFinish. Il primo comando forza tutti i comandi
impilati in una coda ad essere eseguiti sul device a cui è associata la coda ma
non da nessuna garanzia che quando ritorni il controllo al programma i comandi
siano niti.
Per ovviare a questo, la funzione clFinish ha il compito di bloccare il usso
di programma dell'host e di ritornare il controllo se e solo se i comandi impilati
nella coda specicata hanno completamente terminato la loro esecuzione.
funzione clFinish è un punto di blocco per il programma host.
La
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
42
2.2.5 Conclusioni
Questa sezione è servita per mostrare un'idea generale sulla specica del linguaggio OpenCL. Non ha assolutamente lo scopo di sostituire la specica e non
andrebbe presa come reference. Per ogni particolare bisogna riferirsi alla fonte
[24].
Inoltre non è stata arontata tutta la specica OpenCL ma solo una piccola
parte di questa, necessaria a capire l'esempio che seguirà e necessaria a capire
le scelte implementative del framework senza dover per forza padroneggiare e
conoscere tutte le chiamate e le caratteristiche delle funzioni OpenCL.
OpenCL è un linguaggio molto potente e fornisce i mezzi al programmatore
per sfruttare tutta la forza all'interno di una macchina. Occorre ricordare ed
appuntarsi una nota sul fatto che non tutti i problemi possono (e devono) essere
risolti in OpenCL. Una delle fasi più complicate consiste infatti nel capire come
parallelizzare un problema e se questo problema è eettivamente parallelizzabile.
In caso si forzi questo meccanismo, OpenCL si trasforma da un linguaggio
utile ad uno strumento male usato e questo porta, come la maggior parte delle
volte che si agisce secondo questo comportamento, ad un programma ineciente,
di dicile comprensione e, molto probabilmente, errato.
2.3 Un esempio di OpenCL
In questa sezione verrà mostrato un esempio di codice OpenCL, dalla creazione
del programma host e del kernel alla vera e propria esecuzione del codice.
Il programma in esempio somma due array, cella per cella, e scrive il risultato
in un altro array. E' stato scelto questo programma perchè è il classico Hel-
lo World del calcolo GPGPU e perchè è semplice e permette di capire come
funzionano i meccanismi basilari.
Farò riferimento ai tutorial e/o al codice
di[13, 25].
Un programma semplice come questo permette di identicare una struttura generale di un codice OpenCL dividendo il programma in cinque grandi
parti: inizializzazione delle risorse, compilazione del programma, creazione degli
oggetti in memoria, esecuzione dei kernel e lettura dei dati e release delle risorse.
La progettazione del kernel, obiettivo primario del programmatore, viene
discussa come primo argomento.
Il codice in queste pagine ha il solo scopo di dare una visione generale di
come può essere una computazione OpenCL dall'inizializzazione dell'ambiente
alla ne della computazione.
2.3.1 Il kernel
Il design di un kernel dovrebbe essere la prima parte della realizzazione di un
programma OpenCL. Obiettivo del programmatore è quello di ottenere il massimo dell'ecienza dalla macchina su cui sta lavorando o creare codice che possa
funzionare il più ecientemente possibile su un gruppo eterogeneo di macchine.
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
43
Davanti ad un problema il programmatore deve chiedersi: Posso parallelizzare il problema?
La sua natura si presta a dividere l'esecuzioni in più parti
indipendenti? Se si, quali? Quali strumenti mi fornisce OpenCL per ottimizzare
la gestione della memoria?
Spesso si possono ottenere enormi dierenze con scelte implementative leggermente diverse, assolutamente indierenti per quanto riguarda il sorgente, ma
totalmente diverse per quanto riguarda il codice generato per il device.
In questo esempio, questo kernel permette di sommare due array, cella per
cella, e di mettere il risultato nel terzo array. Si può subito notare il parallelismo;
il codice del kernel infatti è molto simile ad un codice C che compie lo stesso
compito, a meno di un ciclo.
Vediamo perchè.
Il codice C che svolge questo compito è costituito da un ciclo che per ogni cella dell'array, somma la i-esima cella del primo array con la i-esima cella del secondo array e mette il risultato nella i-esima cella del terzo array(∀ i [0, count −
1]):
//Codice C per sommare due array cella per cella
void add(
float *primo,
float *secondo,
float *risultato,
int length){
int i;
for(i = 0; i < length; ++i){
risultato[i] = primo[i] + secondo[i];
}
}
mentre il codice del kernel OpenCL è:
//Kernel OpenCL per sommare due array cella per cella
__kernel void add(
__global float *primo,
__global float *secondo,
__global float *risultato){
int i = get_global_id(0);
risultato[i] = primo[i] + secondo[i];
}
Come scritto prima, la grande dierenza è l'assenza, nel codice OpenCL, del
ciclo.
Come si agisce in OpenCL?
CAPITOLO 2.
44
OPENCL: OPEN COMPUTING LANGUAGE
In OpenCL, quando deniamo uno spazio di indici, la chiamata di esecuzione
di un kernel esegue un'istanza di kernel, work-item, su ogni indice presente nello
spazio degli indici.
Mentre il codice C sopra scritto verrà eseguito solo una volta (in cui ci sarà un
ciclo denito da
i=0
alla lunghezza degli array
length − 1),
un kernel OpenCL
viene eseguito tante volte quanti sono gli indici appartenenti allo spazio degli
indici. La funzione get_global_id serve infatti per interrogare il sistema e
sapere quale kernel sono nello spazio di lavoro dei kernel . Questa piccola, ed
allo stesso tempo enorme, dierenza sta alla base di ogni programma OpenCL.
E' quindi compito del programmatore usare questo strumento come meglio può
per risolvere i problemi che gli vengono posti volta per volta.
2.3.2 Inizializzazione delle risorse
Supponendo di avere già gli array allocati nella memoria dell'host, scriviamo solo
la funzione che gestisce l'ambiente OpenCL e le sue chiamate. In ingresso prenderà quindi tre puntatori ad array di, ad esempio,
la lunghezza degli array
f loat
17 . La funzione ritorna un valore
e un n rappresentante
int , CL− SU CCESS ,
se la funzione ha svolto il calcolo correttamente.
//include dei file headers di OpenCL
#include <CL/cl.h>
//stringa del kernel add
const char *stringaProgramma = "\n" \
"__kernel void add(
"__global float *primo,
"__global float *secondo,
"__global float *risultato)
"{
"
int i = get_global_id(0);
"
risultato[i] = primo[i] + secondo[i];
"}
"\n";
/**
\n"
\n"
\n"
\n"
\n"
\n"
\n"
\n"
\
\
\
\
\
\
\
\
Funzione che imposta l'ambiente OpenCL ed esegue
la computazione. Codice eseguito su cpu:
GenuineIntel, Intel(R) Pentium(R) D CPU 2.66GHz
*/
int runCL(
float *primoArray,
float *secondoArray,
17 Per
comodità supponiamo gli array di lunghezza uguale. Lo scopo di questa parte è quella
di dare un'idea dei meccanismi di OpenCL di base senza guardare a controlli o correttezza
formale del codice dell'host.
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
float *risultato,
int lunghezzaArray) {
//oggetto programma
cl_program programma;
//oggetto kernel
cl_kernel kernel;
//coda dei comandi
cl_command_queue coda_comandi;
//contesto OpenCL
cl_context
contesto;
//device cpu
cl_device_id deviceCpu = NULL;
//memorizza gli errori che ritornano
//le chiamate OpenCL
cl_int errore = 0;
//grandezza del buffer
size_t grandezza_buffer;
//piattaforma di lavoro OpenCL
cl_platform_id* piattaforma = NULL;
//oggetti di memoria corrispondenti ai
//parametri di ingresso
cl_mem primoArrayMem, secondoArrayMem, risultatoMem;
//numero di piattaforme disponibili
cl_uint numeroDiPiattaforme;
//stringa per registrare il nome del venditore
//del device
cl_char nomeVenditore[1024] = {0};
//stringa per registrare il nome del device
cl_char nomeDevice[1024] = {0};
//global work size
size_t global_work_size = lunghezzaArray;
// Vedo quante piattaforme ci sono:
//sul sistema testato era solo una
errore = clGetPlatformIDs(0,
NULL,
&numeroDiPiattaforme);
piattaforma =
new cl_platform_id[numeroDiPiattaforme];
// Registro la piattaforma
errore |= clGetPlatformIDs(
numeroDiPiattaforme,
piattaforma,
NULL);
45
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
assert(errore == CL_SUCCESS);
// Scelgo come device la CPU
errore = clGetDeviceIDs(
*piattaforma,
CL_DEVICE_TYPE_CPU, 1,
&deviceCpu,
NULL);
assert(errore == CL_SUCCESS);
assert(deviceCpu);
// Prendo alcune informazioni sul device CPU
errore = clGetDeviceInfo(
deviceCpu,
CL_DEVICE_VENDOR,
sizeof(nomeVenditore),
nomeVenditore,
NULL);
errore |= clGetDeviceInfo(
deviceCpu,
CL_DEVICE_NAME,
sizeof(nomeDevice),
nomeDevice,
NULL);
assert(errore == CL_SUCCESS);
printf("Connessione a %s, %s...\n",
nomeVenditore,
nomeDevice);
// Creazione di un contesto su uno specifico device
contesto = clCreateContext(
0,
1,
&deviceCpu,
NULL,
NULL,
&errore);
assert(errore == CL_SUCCESS);
// Creazione coda comandi sul contesto appena creato
coda_comandi = clCreateCommandQueue(
contesto,
deviceCpu,
0,
NULL);
// continua ...
46
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
47
Questa prima parte è abbastanza chiara e non necessita di spiegazioni. Consiste
esclusivamente in una serie di chiamate a funzioni che inizializzano gli oggetti
OpenCL.
2.3.3 Compilazione
In OpenCL il programma viene creato e poi compilato in due modi: in fase di
creazioni si passa una stringa di caratteri, un array di char, oppure il codice
binario compilato del programma.
In questo caso è stato scelto di passare il
programma come stringa di caratteri. Viene poi creato il kernel dal programma
compilato selezionandolo per nome (il nome della funzione kernel).
Il codice si autocommenta:
// continua da sopra ...
// Creo il programma da una stringa
programma = clCreateProgramWithSource(
contesto,
1,
(const char**)&stringaProgramma,
NULL,
&errore);
assert(errore == CL_SUCCESS);
// Compilazione programma
errore = clBuildProgram(programma, 0, NULL,
NULL, NULL, NULL);
assert(errore == CL_SUCCESS);
// Creo l'oggetto kernel da usare per la computazione
kernel = clCreateKernel(programma, "add", &errore);
//continua ...
2.3.4 Creazione degli oggetti memoria
Dopo aver creato i kernel, un programma ha la necessità di creare oggetti che
siano compatibili con i device OpenCL. Vengono creati i tre array ed i primi
due vengono riempiti con i dati passati alla funzione runCL:
// continua da sopra ...
// Alloco memoria necessaria per i
// buffer della memoria del device OpenCL
grandezza_buffer = sizeof(float) * lunghezzaArray;
// Primo oggetto memoria
primoArrayMem = clCreateBuffer(
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
contesto,
CL_MEM_READ_ONLY,
grandezza_buffer,
NULL,
NULL);
// Scrivo dentro il primo oggetto memoria i dati
// del primo array
errore = clEnqueueWriteBuffer(
coda_comandi,
primoArrayMem,
CL_TRUE,
0,
grandezza_buffer,
(void*)primoArray,
0, NULL, NULL);
// Secondo oggetto memoria
secondoArrayMem = clCreateBuffer(
contesto,
CL_MEM_READ_ONLY,
grandezza_buffer,
NULL,
NULL);
// Scrivo dentro il secondo oggetto memoria
// i dati del secondo array
errore |= clEnqueueWriteBuffer(
coda_comandi,
secondoArrayMem,
CL_TRUE,
0,
grandezza_buffer,
(void*)secondoArray,
0, NULL, NULL);
assert(errore == CL_SUCCESS);
// Oggetto memoria risultato
risultatoMem = clCreateBuffer(
contesto,
CL_MEM_READ_WRITE,
grandezza_buffer,
NULL,
NULL);
// Mi assicuro che tutto venga
//scritto prima di continuare
clFinish(coda_comandi);
// continua ...
48
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
49
2.3.5 Esecuzione dei kernel
In questa parte vengono impostati gli argomenti al kernel creato e successivamente viene denito lo spazio di indici, in questo caso 1, e la sua dimensione
(in questo caso
sizeof (f loat) ∗ lunghezzaArray );
dopo aver completato queste
operazioni, viene lanciato il comando di esecuzione.
Dato che il comando di esecuzione ritorna subito il controllo, la chiamata
clFinish blocca il usso del programma per far si che si continui solo quando
l'esecuzione è eettivamente terminata.
Il codice è:
// continua da sopra ...
// Imposto gli argomenti del kernel
errore = clSetKernelArg(kernel, 0,
sizeof(cl_mem), &primoArrayMem);
errore |= clSetKernelArg(kernel, 1,
sizeof(cl_mem), &secondoArrayMem);
errore |= clSetKernelArg(kernel, 2,
sizeof(cl_mem), &risultatoMem);
assert(errore == CL_SUCCESS);
// Faccio partire l'esecuzione
errore = clEnqueueNDRangeKernel(
coda_comandi,
kernel,
1,
NULL,
&global_work_size,
NULL,
0, NULL, NULL);
assert(errore == CL_SUCCESS);
// Aspetto finchè l'esecuzione non finisce
clFinish(coda_comandi);
// continua ...
2.3.6 Lettura dei dati e Release delle risorse
L'ultimo passo di un programma OpenCL è quello di leggere i dati elaborati
dall'esecuzione del kernel. Inoltre bisogna ricordarsi che la creazione di un ambiente OpenCL occupa memoria e occorre liberare questa memoria. In questa
fase quindi si prelevano i dati e si libera la memoria
18 Il
18 :
fatto di liberare la memoria è stato fatto solo per scopi dimostrativi. Un programma
che necessita di eseguire più computazioni OpenCL non è obbligato a pulire ogni volta la
memoria; basta ricordarsi però che c'è della memoria occupata che, quando non deve essere
più utilizzata, deve essere liberata.
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
50
// continua da sopra ...
// Una volta finita l'esecuzione, leggo
// il risultato e lo metto nell'array risultato
errore = clEnqueueReadBuffer(
coda_comandi,
risultatoMem,
CL_TRUE,
0,
grandezza_buffer,
risultato,
0, NULL, NULL);
assert(errore == CL_SUCCESS);
clFinish(coda_comandi);
//libero la memoria da: oggetti memoria,
//kernels, programmi, code e contesti
clReleaseMemObject(primoArrayMem);
clReleaseMemObject(secondoArrayMem);
clReleaseMemObject(risultatoMem);
//kernel, coda comandi, programma e contesto
clReleaseKernel(kernel);
clReleaseProgram(programma);
clReleaseCommandQueue(coda_comandi);
clReleaseContext(contesto);
delete[] piattaforma;
// Ritorno il valore SUCCESS
return CL_SUCCESS;
}
2.4 Conclusioni
In questo capitolo è stata brevemente analizzata l'architettura di OpenCL: gestione degli oggetti dedicati, programmi OpenCL, lettura, scrittura ed elaborazione dei dati.
Bisogna osservare che OpenCL non è nato come linguaggio risolutivo per
ogni problema. Occorre ripetere che questo framework è uno strumento che deve
essere usato per scopi specici ed assolutamente non general purpose . La fase
iniziale, e forse la più delicata, di un programma OpenCL consiste nell'analisi
di un problema e nello studio delle fattibilità in OpenCL.
Questo framework è stato creato per adempiere a speciche richieste.
Se
usato in maniera impropria, questo strumento rischia di rendere boriosa ed
ineciente un'operazione che sarebbe normalmente eseguibile in un contesto di
elaborazione generica di un calcolatore.
Ho voluto dedicare una parte specica di questo capitolo alla realizzazione di
un semplice programma per permette di capire quali sono i nuclei centrali di un
CAPITOLO 2.
OPENCL: OPEN COMPUTING LANGUAGE
51
applicativo OpenCL; questo permetterà di capire più avanti le motivazioni delle
scelte implementative del framework. Molte di queste scelte, infatti, riguardano
la gestione di code, contesti ed oggetti memoria.
In riferimento a quanto scritto, concludo con un saggio detto:
Se userai sempre il martello, tutti i problemi ti sembreranno
chiodi. (Anonimo)
Capitolo 3
Progettazione del Framework
In questo capitolo, dopo aver dato una nozione generale sull'architettura OpenCL ed aver brevemente mostrato su quale hardware lavora, entreremo nel dettaglio della progettazione di questo framework per l'elaborazione delle immagini.
Questo capitolo, come il lavoro da me svolto, si divide principalmente in due
grandi parti:
ˆ
la prima parte analizza la progettazione, le speciche e le scelte implementative della realizzazione di un framework in grado di astrarre gran
parte del meccanismo che agisce alla base di OpenCL; l'obiettivo è quello di
fornire al programmatore uno strumento che possa essere usato in maniera
agile, funzionale, facile e comunque eciente per creare un ambiente per
compiere calcoli in OpenCL;
ˆ
la seconda parte analizzerà una o più applicazioni di questo framework,
non incentrando l'attenzione sullo studio di algoritmi di image processing,
interesse marginale in questa tesi, quanto più sulle scelte implementative svolte nel framework atte a realizzare uno strumento specico per
la programmazione di algoritmi di elaborazione delle immagini. Occorre
far notare come il vincolo di costruire un framework per l'image processing abbia portato alla programmazione di funzioni accessorie senza però
porre costrizioni sull'utilizzo generale dello strumento in sè. Verrà quindi
mostrato uno dei possibili utilizzi di questo framework sperando che da
questo si possano evidenziare le sue potenzialità e versatilità.
Ciò su cui si vuole incentrare l'attenzione del lettore sono le problematiche
che sono state riscontrare durante la progettazione e la programmazione del
framework di OpenCL; inoltre verrà posto l'accento sulle problematiche che
questo framework introduce e sui vantaggi che porta al programmatore delle
librerie in cui questo framework verrà compreso.
52
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
53
3.1 Introduzione alla progettazione:
framework di OpenCL e librerie IVLLIB
La progettazione di questo framework è orientata, come scritto brevemente sopra, alla costruzione di un wrapper di funzioni OpenCL. La scrittura di questo
framework è nalizzata all'integrazione, il più trasparente possibile, con le librerie IVLLIB di image processing.
laboratorio IVL
Queste librerie sono state sviluppate dal
1 del Dipartimento di Informatica dell'Università Milano Bicoc-
ca, laboratorio per cui è stato svolto lo stage e per cui è stata scritta questa
tesi.
Il wrapper OpenCL è stato concepito come strumento di supporto per quegli
algoritmi della libreria che necessitano e sono predisposti al parallelismo; questo
framework verrà sviluppato nel seguente modo:
ˆ
un'astrazione di molte delle funzioni OpenCL ed una gestione più agile
per permettere al programmatore della libreria di aggiungere algoritmi di
image processing senza la boriosità e le meccaniche sintattiche di OpenCL;
ˆ
l'implementazione di alcuni algoritmi per la fase di testing del framework
e delle sue funzionalità.
Questo capitolo è interamente dedicato alle decisioni che hanno portato alla
progettazione ed al seguente sviluppo del framework spiegando le scelte prese e
la losoa di base del software oggetto della tesi.
3.1.1 Le librerie IVLLIB
Nate come raccolta di classi e utilities fra i colleghi del laboratorio IVL, nel
tempo questo set di funzioni si è trasformato in una libreria completa, stabile
ed eciente per tutti gli usi multimediali. Obiettivo della libreria IVL è quello di
fornire al programmatore uno strumento rapido e facile, ma allo stesso potente,
per poter compiere elaborazioni su immagini fornendo strutture dati, funzioni
e classi in grado di adempiere a svariati compiti.
Realizzata interamente in C++, la libreria è stata progettata in maniera
interamente templata, ed è pensata sia per un uso specico delle strutture dati,
sia per un uso generico. La libreria è concepita per supportare il processing di
immagini astraendo quasi completamente da quelle che sono le basi per acquisire
le immagini da le, visualizzarle a schermo ed eettuare operazioni sui pixel di
queste.
Non bisogna sottovalutare le scelte progettuali di queste librerie che permettono ad un programmatore esperto di accedere direttamente ai dati raw ; la
semplicità d'uso non ha minimamente intaccato la potenza e l'ecienza che il
linguaggio di programmazione ad oggetti C++ metteva a disposizione.
La progettazione del framework di OpenCL si è ispirata molto a questa
losoa, cercando di fornire uno strumento al programmatore che renda l'utilizzo
1 http://www.ivl.disco.unimib.it/
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
54
di questo framework estremamente facilitato; questo però non dovrà mai andare
ad intaccare l'ecienza computazionale e di utilizzo di OpenCL, importante
nella libreria IVL ma ancora più importante quando si parla di calcolo parallelo.
Verranno successivamente mostrate le varie scelte progettuali e le speciche
che hanno reso possibile la realizzazione di questo framework.
3.2 Requisiti e speciche del framework
In questa sezione verranno analizzati i requisiti e le speciche adottate durante
la progettazione e la programmazione di questo framework. Come scritto nella
sezione precendente, le scelte sono basate sulla losoa di base della libreria
IVLLIB in cui questo framework sarà inserito.
3.2.1 Requisiti
Verranno qui elencati i requisiti richiesti per la realizzazione del framework.
Tutti i requisiti sono la conseguenza dell'analisi svolta con i relatori cercando
di analizzare quali potrebbero essere i bisogni del programmatore della libreria,
quali meccaniche di OpenCL potrebbero essere astratte e quali caratteristiche
il framework dovrebbe avere.
I requisiti software del framework di OpenCL sono i seguenti:
ˆ
per quanto riguarda la creazione di un contesto OpenCL, si è deciso
di creare un unico contesto che gestisca tutti i device disponibili nel
computer;
ˆ
programmato in C++, il framework sarà un'astrazione ad oggetti delle
meccaniche funzionali di OpenCL;
ˆ
l'utente avrà a disposizione un punto di accesso per usare le funzioni di
OpenCL che consiste in un oggetto statico richiesto tramite un metodo
apposito; per eettuare questa funzionalità è stato usato il desing pattern
del Singleton;
ˆ
l'utente non deve poter maneggiare i programmi direttamente; l'accesso
deve essere controllato e gestito dal framework;
ˆ
l'utente dovrà essere libero di caricare nel framework un numero non determinato di programmi, ovvero l'utente caricherà in tempo reale i vari
kernel che desidera usare se questi non sono già presenti;
ˆ
la libreria IVLLIB dovrà fornire dei programmi OpenCL che dovranno
essere caricati in fase di inizializzazione se si desiderano usare gli algoritmi
della libreria scritti in OpenCL;
ˆ
la gestione delle strutture di OpenCL verrà mascherata da un layer ad
oggetti che avrà il compito di controllare l'uso della memoria e della sua
gestione;
CAPITOLO 3.
ˆ
PROGETTAZIONE DEL FRAMEWORK
55
l'utente deve comunque avere a disposizione tutte le peculiarità di OpenCL
e il livello di astrazione di questo framework deve tentare di non intralciare,
o se inevitabile di farlo il meno possibile, l'utilizzo del framework originario
di OpenCL;
ˆ
l'utente che usa gli algoritmi della libreria IVLLIB scritti in OpenCL non
dovrà avere nessuna conoscenza specica di OpenCL ed il meccanismo
sottostante dovrà essere completamente trasparente;
ˆ
anche se possibile, l'accesso del programmatore ai device non dovrebbe
avvenire direttamente ma dovrebbe passare per delle politiche di scheduling che assegnano automaticamente i device al programmatore;
ˆ
si richiede che il programmatore della libreria abbia una buona conoscenza
del framework sviluppato e di OpenCL.
3.2.2 Speciche
Questa sezione è dedicata alla spiegazione dettagliata di come si intende sviluppare il framework e verrà descritto come risolvere le problematiche che possono
esserci nella progettazione seguendo i requisiti sopra citati. Inoltre verrà posto
l'accento riguardo agli usi possibili del framework da parte del programmatore
della libreria.
Caratteristica principale di questo wrapper OpenCL è l'astrazione del framework già esistente per permettere a chiunque abbia una buona conoscenza di
OpenCL di programmare in maniera agile ed ordinata, senza perdersi in tutte
le funzioni che OpenCL ore, molte delle quali spesso vengono lasciate con i
valori di default.
Obiettivo primario è quello di orire un'interfaccia più chiara che permetta
di scrivere un codice più facile da manipolare e da gestire tenendo conto i due
obiettivi principali che ci si pone nella progettazione e programmazione di questo
framework:
ˆ
l'interfaccia dovrà comunque consentire un accesso libero e diretto a tutte
le risorse OpenCL senza ostacolare un possile utilizzo di una parte del
framework nativo originale e parte di quello progettato per questa libreria;
ˆ
dato che ogni operazione verrà svolta in OpenCL per motivi di ecienza
computazionale, questo framework dovrà cercare di aggiungere il minimo
overhead indispensabile per quanto riguarda l'elaborazione dei comandi.
A conoscenza di questi due punti cardine, si può incominciare a denire le caratteristiche e le sfaccettature che la progettazione di questo framework ha mostrato
durante tutta la fase di costruzione del progetto, no alla programmazione delle
funzioni di prova della libreria.
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
56
3.2.2.1 Astrazione delle meccaniche di OpenCL
Questo punto deve essere analizzato con attenzione durante la fase di progettazione in quanto il livello di astrazione deciderà inevitabilmente la potenza e
l'ecienza del codice sviluppato. Tutta la parte di inizializzazione del contesto
di OpenCL e di inizializzazione del framework (mappe, vettori, ...) avviene una
sola volta quando viene richiesto, e quindi automaticamente creato, l'oggetto
CL_Enviroment.
In seguito, tutta la gestione di richesta code e generazione
dei kernel passa per l'oggetto CL_Enviroment e il programmatore ha il pieno
utilizzo (con le responsabilità che questo comporta) di tutti questi oggetti e delle
operazioni eettuabili con loro.
Oggetti OpenCL come code, kernel, e oggetti memoria saranno tutti mascherati
da classi o strutture in C++.
3.2.2.2 Oggetti
Il framework è programmato con il linguaggio di programmazione ad oggetti
C++. Si è deciso di sfruttare il paradigma della programmazione ad oggetti per
permettere una manipolazione degli strumenti di OpenCL più vicina alla realtà
progettuale.
Occorre osservare una piccola nota: non viene assolutamente forzato il paradigma ad oggetti del C++ ma viene usato solo quando strettamente necessario e si
cerca di creare oggetti il più leggeri possibile per non aggiungere nessun overhead
che non porti con se usi pratici con ciò che aggiunge.
Il paradigma ad oggetti, quindi, sarà presente sempre e solo quando la sua
applicazione avrà notevoli beneci sulla gestione e pulizia del codice rispetto a
quello originario di OpenCL.
Chiaramente questa scelta porterà ad una serie di limiti per quanto riguarda
la responsabilità della gestione della memoria: con l'introduzione degli oggetti,
il programmatore della libreria deve porre molta attenzione alla loro manipolazione; inoltre, come scritto, il codice è pensato per poter essere usato anche
con le funzioni base di OpenCL quindi sarà comunque possibile accedere ai
dati incapsulati nelle classi.
Tutte queste scelte, insieme ad altre discusse in seguito, portano ad una
presa di coscienza da parte del programmatore per quanto riguarda la gestione
di questi oggetti.
3.2.2.3 Singleton Design Pattern
Il Singleton è uno dei design pattern più comuni, elencato anche nel celebre libro
Design Patterns della Gang of Four [29].
Il Singleton fa parte di quelli che vengono deniti design pattern creazionali,
ovvero quella classe di design pattern che ha lo scopo di controllare la costruzione
e la gestione dell'esistenza di un oggetto. In particolare, il singleton ha lo scopo
di garantire che di una determinata classe venga creata una e una sola istanza,
e di fornire un punto di accesso globale a tale istanza.
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
57
Nel caso del framework, infatti, viene usato il metodo della classe CL_Enviroment
per prendere l'unica istanza del programma.
Esempio:
//codice per avere il puntatore all'unica
//istanza di CL_Enviroment
CL_Enviroment *pointerToInstance =
CL_Enviroment::getInstance();
In C++, il Singleton viene sviluppato con costruttore, copy-constructor, operatore di assegnamento e distruttore privati.
3.2.2.4 Kernel OpenCL
La gestione dei kernel e dei le OpenCL è una di quelle funzionalità del framework che ha occupato una buona parte della progettazione. Il framework mette
a disposizione metodi per caricare stringhe contenenti codice OpenCL o direttamente le con codice OpenCL. Una volta che il le o la stringa vengono vericati,
l'utente decide attraverso le politiche di gestione del framework le modalità di
compilazione dei programmi; quando l'utente vuole un kernel OpenCL fra quelli
compilati non farà altro che interrogare l'istanza dell'ambiente OpenCL che il
framework mette a disposizione indicando il nome del kernel compilato.
Con questa modalità un piccolissimo overhead di strutture sovrastanti il
codice sorgente permette di creare ed interrogare comodamente una base di
dati che contiene tutti i kernel caricati nell'ambiente OpenCL.
La libreria IVLLIB verrà fornita con un meccanismo di inizializzazione che
provvederà a caricare nell'ambiente OpenCL tutti i le in cui sono presenti i kernel sviluppati appositamente per la libreria. In caso il programmatore volesse
aggiungere uno o più algoritmi OpenCL alla libreria, basta semplicemente aggiungere il le con i kernel insieme agli altri e caricare il suo le nella procedura
di inizializzazione di OpenCL della libreria.
3.2.2.5 Politiche di gestione
Un'altra peculiarità del framework dovrà essere quella di automatizzare alcune procedure importanti durante l'esecuzione di una computazione OpenCL.
L'utente, oltre ad avere comunque gli strumenti per controllare direttamente
il framework, avrà anche a disposizione delle politiche comportamentali che
deniscono che scelte deve fare il framework per l'utente e quali sono i casi e le
modalità con qui queste scelte devono essere fatte.
Anche questa caratteristica è nata dall'osservazione di comportamenti modali
nei programmi OpenCL: gran parte del codice si basa tutto sulla stessa logica
di fondo, indipendentemente dallo scopo o dalle modalità dell'elaborazione.
I processi che richiedono automazione sono la richiesta di code su specici device, la decisione del momento della compilazione dei programmi e l'elaborazione
dei comandi sui device.
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
58
Nonostante verrà lasciata l'interpretazione di OpenCL di poter scegliere liberamente questi comportamenti, il framework avrà in sé delle politiche per gestire come queste operazioni verranno eseguite automaticamente quando richiesto
dall'esecuzione del normale usso del programma. Le politiche saranno esclusivamente opzioni di priorità per quanto riguarda le scelte di code, device e
momento in cui uno o più programmi vengono compilati.
3.2.2.6 Gestione degli errori
OpenCL ha un meccanismo di gestione degli errori con codice di ritorno. Ogni
funzione ritorna un codice di errore e in base a questo codice l'utente è libero di
gestire l'errore. Inoltre, a meno che l'errore non sia critico (segmentation fault,
...), OpenCL continua la sua elaborazione senza bloccare il codice al primo
errore.
Dato che molte procedure vengono automatizzate e controllate, e quindi
messe in sicurezza, nel framework si è deciso di gestire gli errori di OpenCL
con il meccanismo delle eccezioni del linguaggio C++. Questo comportamento è costruito per consentire un controllo accurato del codice in modo da interromperne l'esecuzione in caso dovesse succedere qualche evento per cui il
normale scorrere dell'elaborazione potrebbe causare uno o più errori.
Mentre per gli errori che possono accadere durante l'esecuzione dei kernel
non c'è rimedio, la standardizzazione e l'automazione di alcune operazioni di
inizializzazione ha permesso di creare del codice che in linea di massima non
dovrebbe portare ad errori.
Con la frase precedente si vuole soltanto aermare che data l'automazione
dei comportamenti basilari di OpenCL, se il programmatore della libreria agisce
secondo quanto scritto nella documentazione fornita con questo software e con
la specica di OpenCL, non si dovrebbe preoccupare di comportamenti strani
del codice senza che questo venga avvertito; inoltre, se l'ambiente viene usato
come scritto, il programmatore della libreria sarà libero di usare tutte le sue
funzionalità in totale sicurezza.
Come scritto in precedenza, questo framework non vuole in nessun modo
limitare l'uso di OpenCL e quindi il programmatore avrà ancora pieno accesso
ai dati che le classi incapsulano ed a tutte le feature basilari di OpenCL.
Rimane ovvio che per gli errori di programmazione verranno aggiunte assert
nel codice per bloccare comportamenti formalmente errati: puntatori a NULL,
dimensioni dei buer errati, ... .
Questo implica che l'uso del framework deve essere fatto da utenti che hanno già una buona padronanza con OpenCL e desiderano semplicemente un
framework che possa astrarre molti dettagli spesso irrilevanti.
3.2.2.7 Trasparenza diversa per utenti diversi
Esistono tre tipologie di utenti per questo framework:
ˆ
il programmatore della libreria IVLLIB che conosce OpenCL e sfrutta
questo strumento per scrivere algoritmi per la libreria astraendo la parte
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
59
di OpenCL e lasciando all'utilizzatore del metodo la stessa interfaccia delle
normali funzioni della IVLLIB;
ˆ
il programmatore che decide di usare OpenCL per sviluppare una funzione
o un particolare algoritmo e decide di usare il framework per scrivere un
codice più pulito, chiaro e sicuro;
ˆ
l'utente della libreria che non ha assolutamente nessuna conoscenza delle
strutture di OpenCL ma ne conosce le potenzialità e vuole utilizzare gli
algoritmi che sono stati scritti usando questo framework.
2
Per quanto riguarda l'utente programmatore , questo conosce bene OpenCL e
ha a disposizione il codice sorgente del framework e la sua documentazione. Con
questi materiali, il livello di trasparenza che questo utente vede è solo relativo in
quanto per usare bene il framework, oltre a sapere OpenCL, il programmatore
dovrà avere anche una buona conoscenza delle scelte progettuali ed implementative eettuate per permettersi di usare questo strumento con comodità durante
la scrittura di algoritmi complessi e funzioni che usano OpenCL. Questo framework permetterà al programmatore di lavorare concentrandosi esclusivamente
sullo scopo nale senza perdere codice e tempo nella creazione di strutture
OpenCL secondarie al ne per cui sta programmando una o più funzioni.
L'utente che utilizza la libreria IVLLIB, invece, non dovrà avere nessuna
conoscenza specica di OpenCL; quest'ultimo avrà solo la peculiarità di scegliere
se elaborare un algoritmo con le normali funzioni della libreria IVLLIB o se usare lo stesso algoritmo, se presente, in OpenCL. Ovviamente questo algoritmo
sarà stato scritto dal programmatore della libreria e l'utilizzatore userà questa
funzione in maniera totalmente indipendente dalla sua implementazione OpenCL. Per concludere, le funzioni usate da quest'ultimo utente sono praticamente
uguali a quelle usate normalmente a meno del fatto che come motore queste
useranno il framework di OpenCL in oggetto della tesi.
Per capire meglio lo scopo degli utenti la gura 3.1 nella pagina successiva
mostra un diagramma dei casi d'uso.
3.3 Diagrammi di usso e UML
In questa sezione verranno mostrati i diagrammi di usso e il diagramma UML
delle classi per sottolineare la logica generale con cui deve essere sviluppato un
codice che usa questo framework e per avere una comprensione generale di come
è strutturata l'architettura del codice.
Questi strumenti, tipici dell'ingenieria del software, sono stati molto utili
per concepire le prime parti del codice e per avere una consapevolezza generale
della struttura che doveva essere costruita.
2 Per
utente programmatore si intende i primi due casi sopra descritti.
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
60
Figura 3.1: Diagramma dei casi d'uso del framework di OpenCL
3.3.1 Flow Chart
La gura 3.2 nella pagina seguente mostra un ow chart di una possibile procedura che sfrutta il framework OpenCL.
Si può facilmente notare come il usso segua sia il programma prima mostrato, che quelli scritti in seguito. La linearità del codice OpenCL consiste principalmente nella sua caratteristica di essere un codice di inizializzazione , ovvero
un codice che imposta un ambiente di esecuzione e lo inizializza; questo rende
il codice poco dinamico.
Il codice del framework, invece, ha un andamento poco meno lineare di questo
e più che altro è composto da scelte che vengono fatte in seguito ai comandi
invocati dall'utente.
Ho preferito mostrare degli esempi e degli schemi di codice lineare per essere
certo che le meccaniche basilari siano chiare. Questo è stato fatto principalmente
3 riuscire ad avere la visione generale di tutto
perchè in OpenCL è spesso dicile
l'applicativo; infatti, grazie alla reference ottima e ben documentata, è chiaro
cosa svolgono le singole funzioni, mentre è meno ovvio assemblare insieme tutti
i mattoncini per creare un programma dalla struttura solida, stabile e sicura.
3.3.2 Diagramma UML delle classi
Con il diagramma in gura 3.3 nella pagina 62 ho voluto mostrare le classi
principali che sono state concepite per la realizzazione di questo framework.
Obiettivo di questo diagramma è quello di dare un'idea degli oggetti di cui
l'utente può disporre quando usa il framework di OpenCL e di come questi
oggetti interagiscono.
3 Questa
è stata la più grande dicoltà arontata nello sviluppo del framework.
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
61
Figura 3.2: Diagramma di usso di una procedura che sfrutta il framework di
OpenCL
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
62
Figura 3.3: Diagramma UML delle classi del framework di OpenCL
Come dallo standard UML, nella gura 3.3 si può intuire la classe più importante: il CL_Enviroment.
Questa classe rappresenta l'ambiente OpenCL,
ovvero il contesto e l'insieme di impostazioni iniziali che vengono automaticamente scelte dall'ambiente stesso in fase di costruzione. Inotre questo oggetto è
rappresentato, come si vede nella gura 3.3 dal link a se stesso, con un Singleton:
è da questo Singleton che tutti gli oggetti vengono generati.
3.3.3 Note progettuali
Dopo aver deciso di rendere il CL_Enviroment l'ambiente di tutte le computazioni, si è deciso di far dipendere la costruzione e la gestione di alcune
parti di OpenCL direttamente da questo oggetto. Come si può vedere, l'istanza di CL_Enviroment controlla l'accesso diretto, mascherandolo all'utente, di
programmi e device.
Per quanto riguarda gli altri oggetti come code e oggetti memoria, dal diagramma è possibile intuire come questi dipendono, direttamente o non, dall'oggetto CL_Enviroment ma non vengono generati e gestiti da questo. Il loro
uso è a totale discrezione e responsabilità dell'utente, nella stessa maniera con
cui lo sarebbero gli oggetti nativi di OpenCL.
Ci sono due oggetti che occorre analizzare in maniera specica: le CL_Queue
e gli oggetti CL_Mem.
Per quanto riguarda le code, queste vengono generate dall'ambiente sotto richiesta dell'utente e vengono associate al device che l'utente ha scelto, se
specicato, o a quello scelto dalla politica.
Gli oggetti CL_Mem sono stati creati templati per poter disporre di un
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
63
container generico dando la possibilità di creare oggetti di memoria di OpenCL
per qualsiasi dato di cui l'utente dispone.
In aggiunta verranno scritte delle specializzazioni di oggetti CL_Mem per
integrarne l'uso con le strutture dati della libreria IVLLIB come vector e matrix
template.
In ultimo occorre far notare un particolare della gura 3.3 nella pagina precedente: tutti gli oggetti hanno un metodo getData(). Questo metodo permette
di accedere ai dati base di OpenCL per lasciare al programmatore la libertà
di poter usare, nello stesso codice, sia chiamate a funzioni del framework nativo
di OpenCL, che chiamate del framework della libreria IVLLIB.
3.4 Design di Algoritmi di Image Processing con
il framework di OpenCL
In questa sezione, la più corposa del capitolo, verranno mostrati due esempi
di progettazione ed implementazione di algoritmi di image processing da parte
di due potenziali utilizzatori del framework: un programmatore della libreria
IVLLIB che scrive il metodo per la gamma correction ed un programmatore
esterno che vuole fare una sua computazione che non è presente nella IVLLIB e
decide, quindi, di usufruire di questo framework. Entrambe le parti verteranno
sull'analisi che dovrebbe compiere il programmatore nel momento in cui decide
di scrivere del codice, sia che questo venga fatto per la libreria, sia che questo
venga scritto come codice dedicato.
Prima di incominciare occorre scrivere una piccola nota: in questa sezione
non verrà mostrato tanto il codice eettivo da cui prendere la sintatti corretta del
framework in quanto non rilevante per lo scopo che questa sezione si pregge.
Verrà invece scritto dello pseudocodice ponendo eventuali dettagli riferiti ad
OpenCL o al framework quando necessario. Obiettivo di questa sezione è infatti
quello di mostrare il vademeecum a cui dovrebbe attenersi chiunque decida di
programmare con questo framework per elaborare computazioni in un ambiente
OpenCL.
3.4.1 Programmatore della libreria: Gamma Correction
La gamma correction è l'operatore puntuale per eccellenza. Prima entrare nel
dettaglio della gamma correction e di come viene applicata, verrà descritta la
natura degli operatori puntuali per capire l'analisi che dovrebbe svolgere un
programmatore quando decide di usare OpenCL.
Gli operatori puntuali sono quelle operazioni che, in un'immagine, agiscono
su un unico pixel indierentemente dagli altri; gli operatori puntuali hanno
quindi due grandi vantaggi:
ˆ
posso applicare l'operazione sull'immagine e scrivere direttamente sullo
stesso pixel in cui sto lavorando in quanto quest'ultimo non inuenzerà
gli altri. Inoltre il fatto di poter scrivere direttamente sullo stesso pixel su
CAPITOLO 3.
64
PROGETTAZIONE DEL FRAMEWORK
cui sto lavorando implica un grande risparmio di memoria in quanto non
devo avere un buer temporaneo per salvare i calcoli;
ˆ
non devo seguire nessun ordine nell'esecuzione e sono libero di incominciare ovunque nell'immagine.
E' intuitivo e quasi scontato capire la natura parallela di questo problema: se
ogni pixel è indipendente uno dall'altro, per ogni pixel verrà creata un'istanza
di un kernel OpenCL che svolgerà l'operazione necessaria. Questo è un ottimo
esempio per mostrare le capacità di OpenCL.
3.4.1.1 La Gamma Correction
La gamma correction è un operatore che è nato direttamente dai problemi che
avevano le televisioni con tubo catodico nel rappresentare le immagini trasmesse
via cavo. Quando un'immagine veniva trasmessa e riprodotta nel televisore a
tubo catodico, l'immagine presentava i colori parzialmente alterati.
Questo
era causato principalmente dalla sica degli elettroni sparati dal cannone del
televisore per colpire i fosfori dello schermo. L'operatore gamma serve quindi a
mappare un range di valori in un dominio più uniforme di quello originale.
La formula della gamma correction è la seguente:
gamma = c ∗ ey
dove
c
e
y
sono i parametri della gamma correction. Normalmente
uguale ad uno ed è il parametro gamma
y
c
è posto
a variare.
3.4.1.2 Progettazione Kernel
Prima di tutto occorre capire come scrivere il kernel di OpenCL e le dimensioni
del problema. Dato che sto eettuando un'elaborazione su una immagine rappresentata logicamente da una matrice, nonostante questa sia immagazzinata
come buer lineare, il nostro problema è un problema a due dimensioni: righe
e colonne della matrice. Il kernel prenderà quindi in ingresso tre parametri: il
vettore rappresentante l'immagine (supponendo l'immagine come RGB8, ovvero
come unsigned char), il parametro
y
della gamma correction e il numero di
colonne della matrice.
Un modello di kernel per elaborare questo calcolo potrebbe essere:
/*
Questo kernel è solo un modello a cui ci si può
ispirare per realizzare una versione di questa
procedura.
Non dovrebbe essere preso come esempio di sintassi
precisa e corretta di OpenCL.
Questa funzione prende in ingresso il buffer
lineare dell'immagine, le colonne dell'immagine
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
65
e il parametro gamma per calcolarne la gamma
correction
*/
__kernel void gamma(unsigned char *image,
float gamma,
int width){
//quì interrogo il sistema per sapere
//che kernel sono, ovvero la posizione
//del pixel nello spazio del
//problema dato che una
//matrice viene risolta come problema
//a due dimensioni
int x = get_global_id(0);
int y = get_global_id(1);
//quì copio il valore, meglio non operare
//direttamente sulla memoria globale
unsigned char value = *image[x * width + y];
//calcolo il valore
unsigned char returnedValue = powr(value, gamma);
//lo scrivo dentro ogni cella
*image[x * width + y] = returnedValue;
}
Se l'immagine ha le dimensioni di 200x200 pixel, verranno lanciate 40000 istanze
di kernel che elaborano su ogni singolo pixel la gamma correction.
3.4.1.3 Possibile Codice del Programma
In questa sotto-sezione viene sviluppato dello pseudocodice per spiegare come
potrebbe essere progettata e scritta una funzione di OpenCL concepita dal
programmatore della libreria IVLLIB che realizza la gamma correction.
Nel programma si presuppone che la funzione venga chiamata solo dopo aver
4
eseguito l'inizializzazione dell'ambiente OpenCL . Si suppone di fare un metodo
che esegua la gamma correction su una matriche di RGB8, ovvero unsigned char,
e che il kernel, chiamato gamma, sia già stato caricato dalla initCL().
void gammaCorrection(Matrix<RGB8> &matrix, float gamma){
//Per sicurezza tutto il codice è inserito
//all'interno di un blocco try-catch
//-------------------------------try{
4 Chiamando
una funzione di inizializzazione come ad esempio la initCL(), che carica tutti
i programmi in memoria.
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
66
//la init deve essere chiamata
//prima di chiamare questa funzione
//--------------------------------//prendo i puntatore all'unica istanza
//dell'ambiente OpenCL
CL_Enviroment *pointer =
CL_Enviroment.getInstance();
//imposta la politica per scegliere
//la CPU come primo device
pointer->setManagement(CPU_FIRST);
//preleva la coda sul device scelto
//dalla politica
CL_Queue queue = pointer->getQueue();
//crea gli oggetti memoria necessari
//all'elaborazione (tipo, puntatore, lunghezza)
CL_Mem cl_matrix(READ_WRITE, matrix.data(),
matrix.columns() * matrix.rows() * sizeof(RGB8));
CL_Mem fGamma(READ_ONLY, &gamma, sizeof(float));
CL_Mem columns(READ_ONLY, &matrix.columns(),
sizeof(int));
//riempi gli oggetti memoria di OpenCL
queue.writeDataToDevice(cl_matrix);
queue.writeDataToDevice(fGamma);
queue.writeDataToDevice(columns);
//creazione del kernel di OpenCL
CL_Kernel ker = pointer->getKernel(gamma);
//vengono impostati gli argomenti
ker.setArgument(0, cl_matrix);
ker.setArgument(1, gamma);
ker.serArgument(2, columns);
//vengono impostate le dimensioni del
//problema OpenCL
ker.setKernelDimension(TWO);
ker.dimensionOne = matrix.columns();
ker.dimensionTwo = matrix.row();
//esecuzione del kernel
queue.run(ker);
//vengono prelevati i dati
queue.readDataFromDevice(cl_matrix, matrix.data());
}catch(exception &e){
//gestione dell'eccezione
}
}
CAPITOLO 3.
67
PROGETTAZIONE DEL FRAMEWORK
3.4.2 Utilizzatore del Framework di OpenCL:
Filtro di Smoothing
Un programmatore vuole applicare un ltro di smoothing alla sua immagine e
dispone già dei parametri che desidera usare. Si rende conto che nella libreria
non è stato implementato un ltro di smoothing e decide di crearne uno e di
usare il framework OpenCL della libreria IVL per scriverlo.
Anche in questo caso il codice scritto deve servire come esempio per il programmatore che si interfaccia con questo framework e desidera avere un'idea
generale di come dovrebbe progettare il codice.
3.4.2.1 Filtro di Smoothing
Il ltro di smoothing viene utilizzato principalmente per sfocare un'immagine
in base ai pesi assegnati al ltro. Lo smoothing è un'operazione di correlazione
e consiste nel passare sull'immagine, pixel per pixel, il ltro costituito da una
matrice di valori che sono pesati in base al fattore di smoothing che vogliamo
applicare.
Quindi, supponendo un ltro quadrato, ad esempio un ltro
3 ∗ 3,
abbiamo
i seguenti dati:
ˆ
per ogni pixel dell'immagine
valore è
ˆ
n ∗ m,
con
x, y
indici dell'immagine, il suo
f (x, y);
per ogni pixel del ltro
3 ∗ 3,
con
i, j
indici del ltro, il suo valore è
g(i, j).
L'operazione può quindi essere riassunta nella seguente formula:
n X
m X
3 X
3
X
(
f (x, y) ∗ g(i, j))
x=0 y=0 i=0 j=0
Si possono notare subito due caratteristiche di questa operazione:
ˆ
la notevole quantità di calcoli necessari per portare a termine tutta la
computazione;
ˆ
l'obbligo di usare di un buer temporaneo dove salvare i dati in quanto,
se li riscrivessi sulla stessa matrice, corromperei i dati per l'elaborazione
successiva.
Dai punti sopra posso quindi dedurre che questa operazione è indipendente pixel
per pixel grazie all'uso di un buer temporaneo; questo porta direttamente a
scoprire la natura parallela del programma e come si adatta perfettamente alla
logica di OpenCL.
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
68
3.4.2.2 Progettazione Kernel
Anche in questo caso partiamo dalla denizione del kernel e dalle dimensioni
del problema. Il problema è ancora una volta a due dimensioni.
Un prototipo del kernel potrebbe essere il seguente:
/*
Questo kernel è solo un modello a cui ci si può
ispirare per realizzare una versione di questa
procedura. Non dovrebbe essere preso come
esempio di sintassi precisa e corretta di OpenCL.
Questa funzione prende in ingresso il buffer
lineare dell'immagine e le colonne dell'immagine
Il filtro viene scritto come costante del kernel
*/
//filtro costante nella memoria
//__constant del device OpenCL
__constant unsigned char[] filter =
{1, 2, 1, 2, 4, 2, 1, 2, 1};
__kernel void convolution(unsigned char *image,
int width){
//quì interrogo il sistema per sapere
//che kernel sono, ovvero la posizione
//del pixel nello spazio del
//problema dato che una
//matrice viene risolta come problema
//a due dimensioni
//-----------------------------------//il +1 serve ad ignorare la cornice
//Vedi commento nel codice dell'host
int x = get_global_id(0) + 1;
int y = get_global_id(1) + 1;
//quì copio il valore, meglio non operare
//direttamente sulla memoria globale
unsigned char value =
*image[x
*image[x
*image[x
*image[x
*image[x
*image[x
*image[x
*image[x
*image[x
*
*
*
+
+
+
1 * width + y - 1] * filter[0] +
1 * width + y] * filter[1] +
1 * width + y + 1] * filter[2] +
width + y - 1] * filter[3] +
width + y] * filter[4] +
width + y + 1] * filter[5] +
1 * width + y - 1] * filter[6] +
1 * width + y] * filter[7] +
1 * width + y + 1] * filter[8];
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
69
value = value / 16;
*image[x * width + y] = value;
}
3.4.2.3 Possibile Codice del Programma
In questa sotto-sezione si ipotizzano alcune linee di codice che un utente che usa
il framework potrebbe utilizzare per compiere un determinato calcolo. Anche
in questa sezione non viene incentrata l'attenzione sulla sintassi; si preferisce
evidenziare, specialmente in questo caso, la possibilità di miscelare il codice del
framework IVL con il codice e le funzioni di OpenCL. Si ipotizza che il kernel
sopra scritto sia in un le smoothing.cl nella stessa cartella dell'eseguibile.
Il codice, direttamente in un main, potrebbe essere:
//...
//Ho una matrix in cui c'è l'immagine
//su cui fare smoothing e sto supponendo che
//il file con il kernel smoothing chiamato
//smoothing.cl sia nello stesso percorso
//dell'eseguibile
//Per sicurezza tutto il codice è inserito
//all'interno di un blocco try-catch
try{
//Non c'è inizializzazione in quanto sto usando
//il framework senza le librerie IVLLIB
//--------------------------------//prendo i puntatore all'unica istanza
//dell'ambiente OpenCL
CL_Enviroment *pointer =
CL_Enviroment.getInstance();
//prendo direttamente la CPU come OpenCL device
CL_Device dev = pointer->getDevice(CPU);
//apro la coda su un device specifico
CL_Queue queue = pointer->getQueue(dev);
//creo e compilo il programma
id_type program =
pointer->addProgramFromFile(smoothing.cl);
pointer->compileProgram(program);
//creo gli oggetti memoria (matrice RGB8)
CL_Mem cl_matrix(READ_ONLY, matrix.data(),
matrix.columns() * matrix.rows() * sizeof(RGB8));
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
70
CL_Mem columns(READ_ONLY, &matrix.columns(),
sizeof(int));
CL_Mem result<unsigned char>(READ_WRITE, NULL,
matrix.columns() * matrix.rows() * sizeof(RGB8));
//riempio gli oggetti memoria
queue.writeDataToDevice(cl_matrix);
queue.writeDataToDevice(columns);
//creo il kernel
CL_Kernel ker = pointer->getKernel(smoothing);
//imposto gli argomenti del kernel
ker.setArgument(0, cl_matrix);
ker.serArgument(1, columns);
//imposto le dimensioni del problema
Dimension dimensions = TWO;
//lo smooth non viene fatto sulla cornice
ker.dimensionOne = matrix.columns() - 2;
ker.dimensionTwo = matrix.rows() - 2;
ker.setKernelDimension(dimensions);
//esegue il kernel
queue.run(ker);
//creo spazio per legger i dati dell'esecuzione
unsigned char cl_data [1000];
//---//lettura dei dati dell'esecuzione
// - per leggere uso codice OpenCL //chiamata OpenCL
clEnqueueReadBuffer(
queue.getData(), //ritorna la coda OpenCL
result.getData(), //oggetto cl_mem
CL_TRUE,
0,
1000 * sizeof(unsigned char),
cl_data, //puntatore ai dati in cui
//metto il risultato
0, NULL, NULL);
//aspetto la coda scrivendo una clFinish esplicita
clFinish(queue.getData());
}catch(exception &e){
//gestione dell'eccezione
}
//...
CAPITOLO 3.
PROGETTAZIONE DEL FRAMEWORK
71
3.4.3 Conclusioni
Il codice sopra scritto mostra la versatilità e la facilità con cui questo framework
può essere usato. Ancora una volta si ricorda che il codice precendentemente
mostrato è solo un esempio fra i tanti possibili ed è lontano dalla correttezza
sintattica con cui andrebbe scritto; i due precedenti esempi rappresentano, con
molta propabilità, i casi possibili di utilizzo del framework.
Durante tutto il codice verranno adeguatamente gestite le eccezioni, se lanciate, e grazie ai campi che queste hanno, per l'utente risulta molto facile
comprendere cosa non ha funzionato nell'elaborazione.
Inne, occorre notare come la sintassi a oggetti permette una scrittura ed
una conseguente lettura del codice molto più facile ed immediata.
3.5 Possibilità di modica ed evoluzione del progetto
Il framework realizzato non si potrà mai denire nito; è infatti semplice notare
due strade evolutive che questo lavoro porta con sé:
ˆ
algoritmi della IVLLIB sempre più complessi e sempre più specializzati
per i device, con ottimizzazioni sempre migliori e per set di dati sempre
più grandi;
ˆ
modiche e upgrade del framework e della sua struttura per orire al
programmatore della libreria uno strumento dalla malleabilità sempre
maggiore senza inuenzare l'ecienza del codice.
Insieme a questo framework viene fornita una documentazione che ha una duplice funzionalità: da una parte si compone della classica documentazione necessaria all'uso del framework ed alle piccole modiche che si vogliono eettuare,
dall'altra una vera e propria guida riassuntiva di OpenCL che ha lo scopo di
spiegare in maniera molto dettagliata come è stato scritto il codice e quali sono
le peculiarità da migliorare e i difetti da correggere.
Risulta quindi rilevante il lavoro riguardo alla documentazione, sempre importante nella programmazione e progettazione del software e di particolare
importanza in questo framework. Uno degli obiettivi principali con cui è stato
scritto è esattamente quello di creare una documentazione accurata che permetta a chiunque studi OpenCL di poter continuare il lavoro sapendo esattamente
le funzioni e le scelte programmative che hanno dato alla luce il codice di questo
progetto.
Capitolo 4
Conclusioni
Questo capitolo mostrerà quali sono le eettive conclusioni di questa tesi, motiverà il lavoro svolto incentrando la questione sull'utilità pratica del linguaggio
ed alla ne verrà lasciato un paragrafo per i commenti personali e il futuro di
questo prodigioso strumento.
4.1 OpenCL: perchè
Questa sezione vuole rispondere a quell'insieme di domande che possono essere
riassunte in:
Cosa porta di nuovo questo framework e perchè dovrebbe essere utilizzato?
4.1.1 Ecienza Computazionale
Al giorno d'oggi, la potenza di un computer è ancora un limite agli obiettivi
che l'essere umano si è posto.
Solo nella computer graphics esistono algorit-
mi ai quali occorrono ore, se non giorni, per restituire un risultato. Calcoli e
simulazioni di ogni tipo eseguono codice parallelo su macchine costruite appositamente e nonostante questo il lavoro dei ricercatori è frenato dai tempi della
macchina.
OpenCL si pone come risposta a tutte queste chiamate, e risponde con un
framework in grado di soddisfare le richieste di chiunque. Inoltre OpenCL rivaluta le architetture multiprocessore commercializzate negli ultimi anni da un
marketing sempre troppo altisonante e mai veramente utile: OpenCL permette,
infatti, di trasformare un processore Quad Core da un centinaio di dollari in
una postazione di calcolo professionale con costi pari a zero.
OpenCL porta in un'azienda della potenza di calcolo praticamente gratis
in quanto permette la rivalutazione delle architetture su cui l'azienda lavora
scrivendo solo software. Inoltre bisogna ricordare che OpenCL ha dimostrato
risultati a dir poco eccezionali con vari prodotti, come ad esempio schede video,
72
CAPITOLO 4.
CONCLUSIONI
73
che appartengono alla fascia commerciale non professionale e quindi dai prezzi
abbordabili e molto competitivi.
4.1.2 OpenCL è Open
La specica di OpenCL è completamente Open ed è in quanto tale sinonimo
di qualità.
Come già scritto, questa specica è diretta dalla Khronos, leader
del mercato e associazione con anni ed anni di lavori dalla qualità eccellente
ed ineccepibile; inoltre questo progetto è svolto con la partecipazione delle più
grandi major del mercato di CPU, GPU e processori embedded. Questo è indice
di qualità ed a dimostrarlo c'è il fatto che, mentre questa tesi viene scritta, la
specica OpenCL 1.1 è in produzione portando correzioni ed ulteriori vantaggi
computazionali.
Il lavoro che la Khronos ha svolto, e continuerà a svolgere, è conseguenza
diretta di standard di altissima qualità e tutte le aziende che hanno sottoscritto
questa collaborazione hanno alle loro spalle anni e prodotti di ottima fattura e
sicurezza: OpenCL è la conseguenza di tutto questo messo insieme.
4.1.3 Con questo framework, OpenCL è facile
Uno degli obiettivi primari di questo framework è quello di facilitare e rendere
più sicuro l'uso di OpenCL da parte di un programmatore esperto. Come più
volte sottolineato, il framework semplica il lavoro di chiunque abbia la necessità di usare le caratteristiche di OpenCL senza però incombere in chiamate a
funzioni sempre uguali e in controlli di errore che si ripetono nel codice.
Ancora una volta occorre sottolineare che questo framework non ha assolutamente la pretesa di far usare OpenCL a qualcuno che non ne conosce il funzionamento; infatti, nonostante una documentazione ricca di contenuto, questo
framework è consigliato solo a coloro che, avendo una buona conoscenza di
OpenCL, possono accedere alla moltitudine di opzioni e utilizzi a cui questo
strumento si presta, sfruttandone ogni linea di codice e liberandone così tutte
le potenzialità.
4.2 Cosa porta di nuovo
Le conclusioni a cui giunge questa tesi sono ormai ovvie: si è riuscito a costruire
un wrapper di funzioni OpenCL e l'obiettivo di fornire un framework che astragga molte funzioni di OpenCL è stato raggiunto. Ora parte del lavoro consisterà
nel continuo miglioramento del framework che dovrà seguire contemporaneamente due vie.
La prima sarà quella di inglobare e rendere automatiche sempre più funzionalità di OpenCL, seguendo le nuove speciche e le nuove caratteristiche che
verranno messe a disposizione per dare al programmatore uno strumento sempre
più intuitivo e sempre più sicuro per programmare con OpenCL.
CAPITOLO 4.
CONCLUSIONI
74
La seconda via consiste nella costruzione di una parte della libreria IVLLIB
in grado di sfruttare questo framework e di dare all'utente, ignaro di come si usi
OpenCL ma conscio delle sue potenzialità, un'interfaccia trasparente da poter
utilizzare per poter eseguire algoritmi sempre più complessi computazionalmente
o per avere un range di funzioni da poter usare sempre maggiore in grado di
fornire un'interfaccia di calcolo che si adatti alle sue esigenze.
4.3 Note nali
L'evoluzione di questo framework, come scritto precedentemente, corre su due
rispettive vie: libreria IVLLIB sempre più ricca di funzioni che usano OpenCL
ed un framework sempre più intuitivo e sicuro da usare.
Le conclusioni di questa tesi arrivano quindi ad una ne:
OpenCL si è dimostrato un valido strumento, pratico, ben progettato e supportato da una implementazione e documentazione che raggiungono i più alti
livelli di eccellenza.
Per l'ultima volta bisogna ricordare l'attenta analisi che occorre fare prima
di sviluppare un'applicazione in OpenCL: questo framework non è una scatola magica che risolve i problemi dei programmatori, ma è un framework che
specializza un ambito della programmazione, fornendo uno strumento in grado
di andare ad un livello inferiore alla normale programmazione procedurale e
strutturale a cui siamo abituati.
Sono questi strumenti che, combinati opportunamente, un giorno renderanno
il computer uno strumento in grado di essere veramente general purpose e di
fornire molte delle risposte che l'essere umano sta ancora cercando e, prima o
poi, troverà.
Bibliograa
http://en.wikipedia.org/wiki/Video_card
[1] Schede Video, Wikipedia,
(ultima visita 2 Maggio 2010)
[2] Tecnologia
FireStream
[3] Tecnologia
AMD
Fire
http://en.wikipedia.org/wiki/AMD_
Stream,
(ultima visita 2 Maggio 2010)
CUDA,
http://en.wikipedia.org/wiki/CUDA
Wikipedia,
(ultima visita 2 Maggio 2010)
[4] AMD,
history
Wikipedia,
http://en.wikipedia.org/wiki/Amd#Corporate_
(ultima visita 2 Maggio 2010)
[5] Storia Tecnologia OpenCL, Wikipedia,
Opencl#History
[6] Serie
GeForce
http://en.wikipedia.org/wiki/
(ultima visita 2 Maggio 2010)
GT200,
GeForce_200_Series
Wikipedia,
http://en.wikipedia.org/wiki/
(ultima visita 2 Maggio 2010)
[7] NVIDIA GTX285, sito proprietario NVIDIA,
object/product_geforce_gtx_285_it.html
http://www.nvidia.it/
(ultima
visita
2
Maggio
2010)
[8] GPU,
[9] ATI
http://en.wikipedia.org/wiki/Graphics_
Wikipedia,
processing_unit
5850,
(ultima visita 2 Maggio 2010)
sito
proprietario
manolodeagostini/2009/10/
AMD/ATI,
http://tomsblog.it/
(ultima visita 2 Maggio 2010)
[10] NVIDIA GTX285, sito proprietario NVIDIA,http://www.nvidia.com/
object/product_geforce_gtx_285_us.html
(ultima
visita
2
Maggio
2010)
[11] Descrizione Pipeline di rendering, Wikipedia,
wiki/Graphics_pipeline
http://en.wikipedia.org/
(ultima visita 2 Maggio 2010)
[12] David Luebke, Nvidia Researcher e Greg Humphreys, University of Virginia, How GPUs Work,
paper.php?paper_id=59
http://www.cs.virginia.edu/~gfx/papers/
(ultima visita 2 Maggio 2010)
75
76
BIBLIOGRAFIA
[13] David Gohara, Center of Computational Biology, Responsabile dei tutorial di OpenCL,
http://www.macresearch.org/opencl
(ultima visita 2
Maggio 2010)
[14] Fonte
OpenCL
Tutorials,
sito,
http://www.macresearch.org/Opencl
(ultima visita 2 Maggio 2010)
chip
GT200,
http://www.beyond3d.com/images/reviews/
gt200-arch/gt200die-big.jpg (ultima visita 2 Maggio 2010)
[15] Fonte
chip
GT200,
http://www.xbitlabs.com/images/video/
geforce-gtx200-theory/g200chip.png (ultima visita 2 Maggio 2010)
[16] Fonte
chip GT200, http://www.bit-tech.net/hardware/graphics/
2008/06/24/nvidia-geforce-gtx-280-architecture-review/7
[17] Layout
(ultima visita 2 Maggio 2010)
[18] Fonte chip Core 2 Duo,
core2.03.jpg
[19] Shader
Shader
[20] Frame
Buffer
[21] ROP,
http://www.sudhian.com/img/intel/core2/
(ultima visita 2 Maggio 2010)
programmabili,
Wikipedia,
http://en.wikipedia.org/wiki/
(ultima visita 2 Maggio 2010)
Buer,
Wikipedia,
http://en.wikipedia.org/wiki/Frame_
(ultima visita 2 Maggio 2010)
http://en.wikipedia.org/wiki/Render_Output_
Wikipedia,
Processor
(ultima visita 2 Maggio 2010)
[22] Recensione
Architettura
http://www.
GTX280,
bit-tech.net/hardware/graphics/2008/06/24/
nvidia-geforce-gtx-280-architecture-review/4
(ultima
visita
2
Maggio 2010)
http://www.hwstation.net/
recensioni/evga_geforce_gtx_280_sc/gt200-600.html (ultima visita
[23] Figura e recensione Architettura GTX280,
2 Maggio 2010)
[24] Khronos, OpenCL specication,
http://www.khronos.org/registry/cl/
(ultima visita 2 Maggio 2010)
[25] Justin Hensley, Ph.D. Senior Member of Technical Sta, Oce of the
CTO
-
Advanced
Technology
Initiatives
AMD/ATI,
Video
Tutorials
http://developer.amd.com/documentation/videos/
OpenCLTechnicalOverviewVideoSeries/Pages/default.aspx (ultima
on
OpenCL,
visita 2 Maggio 2010)
[26] Single Instruction Multiple Data, Wikipedia,
wiki/SIMD
(ultima visita 2 Maggio 2010)
http://en.wikipedia.org/
77
BIBLIOGRAFIA
[27] Single Process Multiple Data, Wikipedia,
wiki/SPMD
http://en.wikipedia.org/
(ultima visita 2 Maggio 2010)
[28] Architettura di memoria di un device OpenCL,
p/zocle/wiki/MemoryModel
[29] Singleton Design Pattern,
http://code.google.com/
(ultima visita 2 Maggio 2010)
http://it.wikipedia.org/wiki/Singleton
(ultima visita 2 Maggio 2010)
[30] Utilizzo di CUDA nell'implementazione di algoritmi di supporto all'Elaborazione delle Immagini , Tesi di laureat triennale di Pigazzini Andrea,
2007/2008