Transcript Soutenance

Traitement de l’information
Accélération d’algorithmes de
convolution sur processeurs graphiques
Programmation massivement parallèle
Benoît Pencrec'h
Arrière plan: Nvidia.fr
Accélération d’algorithmes de convolution sur processeurs graphiques
1 - Convolution matricielle et processeur graphique : un couple idéal
1.1 - Convolution et complexité arithmétique
1.2 - Les processeurs graphiques, une architecture taillée pour le calcul
1.3 - Un couple permettant une exécution hautement parallèle
2 - Implémentations et optimisations
2.1 - Programmer sur GPU avec l’API Cuda™
2.2 - Convolution 2D
2.3 - Convolution 2D optimisée
2.4 - Convolution 3D, une histoire de boucle…
3 - Résultats : qualité, performances et limites
3.1 - Qualité des résultats
3.2 - Benchmark des convolutions 2D et 3D
3.3 - Les transferts mémoires, goulot d’étranglement
3.4 - Un autre moyen d’exécuter la convolution : Fourrier
13/04/2015
Benoît Pencrec'h
2
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Convolution et complexité arithmétique
• La convolution mesure la quantité de chevauchement entre deux fonctions
• Pour des termes de type discret elle s’exprime
• La convolution d’un pixel N d’une image I par un noyau K de rayon R s’exprime :
13/04/2015
Benoît Pencrec'h
3
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Convolution et complexité arithmétique
Noyau
Pixel source
Pixel résultat
13/04/2015
Benoît Pencrec'h
Schéma : université Harvay Mudd
4
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Convolution et complexité arithmétique
• Pour une image de taille m . m et un noyau de taille n . n, la convolution
demande n² . m² multiplications
30 s pour convoluer une image 8192.8192 Px
(noyau 9.9)
• En 3D, la convolution requiert n³. m³ multiplications !
7 min 30 s pour convoluer un volume 1024.1024.1024 Px
(noyau 3.3.3)
13/04/2015
Benoît Pencrec'h
5
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Résultats : performances et
limites
Implémentations et optimisations
Convolution et complexité arithmétique
• Il existe une solution pour diminuer le nombre de calculs :
la séparation du noyau
*
abaisse le nombre de multiplications à 2n.m² en 2D et 3n.m³ en 3D
13/04/2015
Benoît Pencrec'h
Schémas : réalisation personnelle
6
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Les processeurs graphiques, une architecture taillée pour le calcul
• Habituellement destinés aux rendus de scènes 3D (jeux vidéos, CAO, …)
• Autrefois spécialisé dans la 3D, ils est devenu universel
• Gère moins d’instructions que le CPU mais possède une puissance de
calcul brute bien plus importante
13/04/2015
Benoît Pencrec'h
7
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Résultats : performances et
limites
Implémentations et optimisations
Les processeurs graphiques, une architecture taillée pour le calcul
Puissance de calcul des GPU Nvidia vs CPU Intel
13/04/2015
Benoît Pencrec'h
Graphe : nvidia programming guide
8
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Les processeurs graphiques, une architecture taillée pour le calcul
CPU
13/04/2015
GPU
Benoît Pencrec'h
Schémas : nvidia programming guide
9
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Les processeurs graphiques, une architecture taillée pour le calcul
• 1 thread
Stream
Processor
• Exécution séquentielle
• 2 ko de registres
13/04/2015
Benoît Pencrec'h
Schémas : nvidia programming guide
10
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Les processeurs graphiques, une architecture taillée pour le calcul
• 8 threads (bloc)
SIMT
(Single Instruction Multiple Threads)
• Exécution parallèle
• 16 ko de mémoire partagée
13/04/2015
Benoît Pencrec'h
Schémas : nvidia programming guide
11
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Les processeurs graphiques, une architecture taillée pour le calcul
• plusieurs SIMT
• exécution parallèle
• mémoire globale
GPU Complet
13/04/2015
Benoît Pencrec'h
Schémas : nvidia programming guide
12
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Un couple permettant une exécution hautement parallèle
• Pour exprimer leur puissance, les GPU requièrent des calculs
indépendants
• La convolution possède cette particularité :
chaque pixel résultat est une somme de
produits indépendante du résultat voisin
• L’utilisation de la convolution sur processeurs graphiques est donc
pertinente
13/04/2015
Benoît Pencrec'h
13
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Programmer sur GPU avec l’API Cuda™
• CUDA™
Compute Unified Device Architecture
• API propriétaire apparue en 2006 avec l’architecture G80 (Ge force 8)
• Se présente comme une extension du langage C (compilateur
spécifique NVCC)
• Permet d’écrire des fonctions C dénommées Kernels qui seront
exécutées N fois en parallèle sur N différents threads
13/04/2015
Benoît Pencrec'h
14
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Programmer sur GPU avec l’API Cuda™
Addition de vecteurs en Cuda
Kernel GPU
Parallèle
// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
int i = threadIdx.x;
C[i] = A[i] + B[i];
}
Code CPU
Séquentiel
int main()
{
...
// Kernel invocation with N threads
VecAdd<<<1, N>>>(A, B, C);
}
• Chaque thread va effectuer une des additions
• N additions seront exécutées simultanément
13/04/2015
Benoît Pencrec'h
Extraits de code : nvidia programming guide
15
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Programmer sur GPU avec l’API Cuda™
• Pour simplifier l’exécution de tâches plus complexes, les threads
bénéficient d’une hiérarchie bien définie.
• L’indice de thread (threadId) peux prendre une (thradId.x) deux
(threadId.y) ou trois (threadId.z) dimensions.
• Apparaissent dès lors blocs de threads (blockId) à une deux ou
trois dimension suivant la même structure.
• L’ensemble des blocs forment une grille de thread (1D ou 2D mais
pas 3D) qui est ensuite envoyée au GPU
13/04/2015
Benoît Pencrec'h
16
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Programmer sur GPU avec l’API Cuda™
Représentation de la grille de thread
13/04/2015
Benoît Pencrec'h
Schémas : nvidia programming guide
17
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Programmer sur GPU avec l’API Cuda™
Addition de matrices en Cuda
Chaque thread du bloc
exécute une addition
Appel du kernel avec 1
bloc de N*N threads
13/04/2015
// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
int i = threadIdx.x;
int j = threadIdx.y;
C[i][j] = A[i][j] + B[i][j];
}
int main()
{
...
int numBlocks = 1;
dim3 threadsPerBlock(N, N);
MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
}
Benoît Pencrec'h
Extraits de code : nvidia programming guide
18
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Convolution 2D
• On découpe l’image en blocs de 16*16 threads
• Chaque pixel résultat est calculé par un thread
• Un thread exécute séquentiellement les sommes / produits nécessaires
au calcul d’un pixel résultat
• Tous les pixels résultats sont calculés en parallèle
13/04/2015
Benoît Pencrec'h
19
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Convolution 2D
• Il reste cependant un problème : les effets de bord
13/04/2015
Benoît Pencrec'h
Schémas : réalisation
personnelle
20
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Convolution 2D
• solution : placer des zéros autour de l’image
• Dans le kernel on choisira de faire des tests sur les indices de thread (plus
performant)
13/04/2015
Benoît Pencrec'h
Schémas : réalisation
personnelle
21
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Convolution 2D optimisée
• Cette implémentation est basée sur l’exemple ConvolutionSeparable du
SDK Nvidia™
• Première optimisation : utilisation de la convolution séparable
l’exécution est divisée en deux passes (i.e. deux kernels)
Une passe horizontale qui traite des rangées de 16 pixels
Une passe verticale traite des colonnes 16 pixels
13/04/2015
Benoît Pencrec'h
22
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Résultats : performances et
limites
Implémentations et optimisations
Convolution 2D optimisée
• Deuxième optimisation : chargement de l’image directement dans la
mémoire partagée du GPU (mémoire shared)
Sans mémoire shared
13/04/2015
Avec mémoire shared
Benoît Pencrec'h
Schémas : nvidia programming guide
23
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Convolution 3D, une histoire de boucle…
Problème: On
: ondécoupe
ne peuxlepas
définirenune
grille
de thread
en 3D
•• Solution
volume
slices
traitées
par nos
kernels 2D
y
x
Z slices
z
Z slices
13/04/2015
Benoît Pencrec'h
Schémas : réalisation
personnelle
24
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Convolution 3D, une histoire de boucle…
• Cette méthode nécessite cependant d’adapter les kernels 2D
Le noyau étant 3D, les kernels doivent récupérer les
valeurs des slices les entourant
13/04/2015
Benoît Pencrec'h
25
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Résultats : performances et
limites
Implémentations et optimisations
Convolution 3D, une histoire de boucle…
Slice volume n-1
Slice volume n
Slice volume n+1
5
Slice volume n
10
noyau
n-1
Z
5
noyau
n
20
noyau
n+1
13/04/2015
Benoît Pencrec'h
Illustration : réalisation
personnelle
26
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Qualité des résultats
La qualité (i.e la justesse) des résultats est constamment vérifiée par :
• Le calcul de la norme relative entre l’implémentation CPU et la
version GPU (contrôle automatique)
• L’export des images et volumes résultats dans un fichier (contrôle visuel)
Image d’origine
13/04/2015
Image covoluée (Gauss) CPU
Benoît Pencrec'h
Image covoluée (Gauss) GPU
Illustration: réalisation
personnelle
27
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Benchmark des convolutions 2D et 3D
Système de test
13/04/2015
Fabriquant et
modèle
Nombre de
cœurs
Fréquence
des cœurs
CPU
Intel Xeon L5320
4
2.5 GHz
GPU
Nvidia Geforce
GTX295
280
1.24 GHz
Benoît Pencrec'h
28
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Benchmark des convolutions 2D et 3D
Facteur d'acclélération global des différentes convolutions en fonction de la taille de
l'image convoluée
Facteur d'accélération (x vitesse CPU Matlab)
CPU
GPU
GPU Séparable
GPU Séparable avec Shared memory
26.0
20.1
21.0
18.8
18.7
16.5
16.0
x2,5
11.6
9.9
11.0
7.0
6.0
2.1
8.1
8.8
8.0
5.9
3.6
3.3
1.0
512
1024
2048
3072
4096
8192
Résolution (Px)
13/04/2015
Benoît Pencrec'h
29
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Benchmark des convolutions 2D et 3D
Facteur d'acclélération global des différentes convolutions en fonction de
la taille du volume convolué
CPU
GPU
GPU Séparable
GPU Séparable avec Shared memory
56.7
Facteur d'accélération (x vitesse CPU Matlab)
61.0
51.0
44.0
41.0
x3
31.0
20.2
11.0
19.0
16.7
21.0
5.9
13.3
12.6
8.5
6.3
256
384
10.0
1.0
64
128
704
Résolution (Vx)
13/04/2015
Benoît Pencrec'h
30
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Les transferts mémoires, goulot d’étranglement
Facteur d'accélération pour la partie calculatoire des différentes convolutions en
fonction de la taille de l'image convoluée
Facteur d'accélération (x vitesse CPU Matlab)
CPU
GPU
GPU Séparable
GPU Séparable avec Shared memory
1000.0
100.0
307.1
350.1
23.8
28.0
353.4
359.7
62.4
44.5
14.0
11.8
10.0
6.9
3.6
x65 !
15.8
8.4
5.6
4.4
5.5
2.7
1.0
512
1024
2048
3072
4096
8192
Resolution (Px)
13/04/2015
Benoît Pencrec'h
31
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Implémentations et optimisations
Résultats : performances et
limites
Les transferts mémoires, goulot d’étranglement
Pourcentage de transfert mémoire par rapport au temps d'éxecution total en
fonction de la résolution de l'image convoluée
GPU
GPU Séparable
GPU Séparable avec Shared memory
100
Pourcentage de transfert mémoire
90
95.3
94.8
94.6
94.6
58.3
58.7
41.1
42.3
94.7
94.4
80
70
75.6
60
62.1
50
40
61.7
44
30
44.5
29.7
20
29.8
18.3
10
0
512
1024
2048
3072
4096
8192
Résolution (Px)
13/04/2015
Benoît Pencrec'h
32
Accélération d’algorithmes de convolution sur processeurs graphiques
Convolution matricielle et
processeur graphique : un couple
idéal
Résultats : performances et
limites
Implémentations et optimisations
Un autre moyen d’exécuter la convolution : Fourrier
Vitesse d'éxecution totale des convolutions 2D GPU linéaires et fréquentielles en
fonction de la taille du noyau
Convolution 2D GPU Séparable avec Shared
Convolution 2D FFT GPU
160
140
Vitesse (MPx/s)
120
100
80
60
40
20
0
1
2
4
8
16
32
64
128
256
512
Taille noyau (Px)
13/04/2015
Benoît Pencrec'h
33
Accélération d’algorithmes de convolution sur processeurs graphiques
Une expérience enrichissante
•Cuda est un langage simple à mettre en œuvre donnant de bons résultats
•Au cours de ce stage j’aurai:
•Acquis une bonne connaissance des GPU et de la programmation parallèle,
•Appris à interfacer des fonctions C et CUDA avec Matlab,
•Utilisé un outil de versionning de Code,
•Familiarisé avec un OS libre,
•Appris à développer sans IDE en utilisant l’outil Make et les makefiles,
•Approfondi mes connaissances en langage C.
13/04/2015
Benoît Pencrec'h
34
Accélération d’algorithmes de convolution sur processeurs graphiques
Sources
• Index de la documentation Nvidia CUDA
Site : http://developer.nvidia.com/
•Université Harvay Mudd
Site : http://fourier.eng.hmc.edu/
13/04/2015
Benoît Pencrec'h
35