Rapport de Stage

Rapport de Stage

-

33 pages
Lire
Le téléchargement nécessite un accès à la bibliothèque YouScribe
Tout savoir sur nos offres

Description

  • rapport de stage
Rapport de Stage « Étude pour l'accélération du code de calcul parallèle elsA à l'aide de   processeurs graphiques (GPU). » Université Pierre et Marie Curie Paris VI Master de Mathématiques et Applications Spécialité : Mathématiques de la Modélisation Parcours: Mathématiques & Informatique Année 2008/2009 Florent Dahm 1
  • texture cudabindtexturetoarray
  • lecture de la texture
  •  l'appel
  • simulation   pour   des   secteurs
  • phase   du   travail
  • exécution   de   threads
  •  comparaison des puissances de calcul des cpu et des gpu ¶
  •  grille  
  • systèmes  
  • cadre du stage

Sujets

Informations

Publié par
Nombre de lectures 2 409
Poids de l'ouvrage 2 Mo
Signaler un problème

Rapport de Stage
« Étude pour l’accélération du code de calcul parallèle elsA à l’aide de 
processeurs graphiques (GPU). »
Université Pierre et Marie Curie
Paris VI
Master de Mathématiques et Applications
Spécialité : Mathématiques de la Modélisation
Parcours: Mathématiques & Informatique
Année 2008/2009
Florent Dahm
1Table des matières
I. Cadre du stage.................................................................................................................................................. ..........3
1. L'entreprise...................................................................................................................................... .....................3
2. Contexte du stage...................................................................................................................... ...........................4
3. Mission et déroulement......................................................................................................................... ...............4
II. Nvidia Cuda et le GPU T10..................................................................................................................................... .5
1. Architecture.................................................................................................................................. ........................5
a) Modèle d'organisation des threads................................................................................................... ...............6
b) Organisation hiérarchique de la mémoire............................................................................................ ...........7
2. Modèle de programmation..................................................................................................................... ..............9
a) Kernels et threads.............................................................................................................................. ..............9
b) Utilisation de la mémoire.................................................................................................................. ............10
c) Compilation, exécution et performance................................................................................... .....................11
3. Premiers essais.................................................................................................................................................... 13
III. Analyse du code elsA.......................................................................................................................... ..................14
1. Présentation du code............................................................................................................................... ............14
a) Historique ............................14
b) Architecture générale.............................................................................................................. ......................15
c) Cas tests utilisés......................................................................................................................................... ....15
2. Profiling ............................................................................................................................................ .................17
3. Analyse du profiling et sélection des noyaux à porter.............................................................................. .........20
IV. Portage sur GPU des routines et résultats......................................................................................... ....................21
1. Calcul de distance................................................................................................................................. ..............21
a) Schéma algorithmique......................................................................................................................... ..........21
b) Validation du calcul et premiers résultats........................................................................................... ..........22
c) Optimisation et relevé de performance.......................................................................................... ...............23
2. Calcul de flux.................................................................................................................................... .................25
a) Schéma algorithmique................................................................................................................................ ...25
b) Vérifications des calculs et premières mesures.................................................................................. ..........25
c) Améliorations et analyse des performances.............................................................................................. ....26
V. Conclusion.......................................................................................................................................... ....................29
 ANNEXE A: Rapport d'installation............................................................................. ..30
 ANNEXE B: Modifications apportées au code elsA....................................................32
 BIBLIOGRAPHIE...................................................................................... ...................33
2I. Cadre du stage
1. L'entreprise
J'ai été accueilli du 9 mars au10 juillet 2009 au sein de département « Systèmes Haute Performance » de CS, sous la 
tutelle de François Roudot.
CS, Communication & Systèmes, est une SSII qui travaille essentiellement sur des projets sensibles et complexes.
CS s'appuie sur ses centres d'excellence technologique pour répondre aux enjeux de ses clients : Simulation et réalité 
virtuelle,   systèmes   embarqués,   sécurité   des   systèmes   d'information,   sûreté   et   continuité   de   fonctionnement,   systèmes 
d'information techniques (PLM, GED, SIG…), qualification logicielle, logiciels libres, génie logiciel et système... 
Ces centres d'excellence technologique permettent à CS d'apporter à ses clients des systèmes clés en main innovants 
et d'assurer la performance et la pérennité des systèmes et des infrastructures critiques.
Intégrateur de solutions clés en main, CS poursuit activement le développement de produits matures à forte valeur 
ajoutée, pouvant facilement s'intégrer dans des solutions complètes. 
Illustration 1: Chiffres clés CS
La simulation fait partie des technologies clé éminemment transverses. Aussi, CS développe des programmes de 
simulation   pour   des   secteurs   très   variés   comme   l’aéronautique,   le   spatial,   l’énergie,   la   défense,   l’automobile,   la 
météorologie, la gestion des risques liés aux catastrophes naturelles, terroristes, …
Dans ce contexte, le calcul haute performance a pour but la résolution de problèmes industriels colossaux en 
s'appuyant   sur   des   technologies   de   plus   en   plus   avancées,   tant   du   point   de   vue   théoriques   avec   des   modélisations 
mathématiques fines, que de celui des moyens informatiques mis en œuvre, impliquant des clusters de calculs comportant 
des milliers de processeurs fonctionnant en parallèle.
32. Contexte du stage
Le code de calcul elsA (ensemble logiciel de simulation en Aérodynamique) est issu d’un programme de l’ONERA 
ayant pour objectif de proposer un environnement unique pour la simulation numérique de la dynamique des fluides.
Le défi actuel est de lancer des simulations se rapprochant toujours plus de la réalité, ce qui induit des volumes de 
données à traiter toujours plus importants. Cette montée en puissance passe par une augmentation forte du temps de 
restitution. Pour réduire ce temps, plusieurs stratégies peuvent être envisagées :
- la multiplication du nombre de processeurs impliqués dans le calcul
- la réduction du temps d’exécution séquentielle par l’utilisation d’accélérateurs adossés au CPU.
Ces derniers mois, les centres de calcul de pointe (CCRT, CINES, ...) se sont équipés en accélérateurs graphiques 
pour améliorer les performances des codes. En parallèle, la couche logicielle pour programmer ces GPU arrive à un niveau 
de maturité permettant de l’intégrer dans des codes industriels.
3. Mission et déroulement
L'objet du stage proposé était de faire une étude du code elsA pour en extraire les noyaux de calcul pouvant faire 
l’objet d’un portage sur les processeurs graphiques.
Le stage s'est déroulé en 3 parties.
Une   première   phase   du   travail   a   consisté   en   une  étude   des   architectures   et  des   technologies   logicielles   des 
processeurs graphiques afin d’appréhender les problématiques du calcul sur GPU.
Dans un second temps, après une période de prise en main du code elsA et de son architecture, celui ci a été analysé 
afin d’en extraire les zones de codes propices à un portage sur GPU.
La troisième étape a consisté au portage des noyaux de calculs sélectionnés avec la librairie de programmation 
CUDA spécifique aux cartes graphiques NVIDIA.
4II. Nvidia Cuda et le GPU T10
Bien qu'initialement conçu pour le rendu graphique, les processeurs graphiques (GPU) d'aujourd'hui peuvent être vu 
comme des processeurs parallèles dotés d'une couche logicielle qui permet de les programmer et d'exploiter leur puissance 
de calcul.
Ce chapitre présente brièvement le calcul sur GPU ou GPGPU (pour « General­Purpose Computation on Graphics 
Processing Units »  en anglais) en s'appuyant sur le processeur graphique Nvidia T10, présent dans les cartes de calculs 
Nvidia Tesla C1060, et sur l'environnement de développement Nvidia Cuda.
Bien que Cuda (pour « Compute Unified Device Architecture ») représentait initialement l'architecture des GPU 
Nvidia programmables, le nom est généralement également associé à celui de l'api «C pour Cuda». Dans ce document, le 
terme Cuda désignera aussi bien l'architecture que la surcouche de programmation.
Nota bene: le GPU T10 est une extension du GT200, les coeurs des processeurs tournent à une fréquence un peu 
plus élevée et la puce embarque plus de mémoire mais l'architecture reste la même.
1. Architecture
Depuis 2003, l'augmentation des fréquences des CPU a drastiquement diminué à cause de trop forte consommation 
d'énergie. Les fabricants de processeurs ont donc fabriqué de plus en plus de modèles multi­cœurs.
C'est à partir de ce moment là que la puissance de calcul des GPUs a commencé à devancer celle des CPUs, comme 
le montre le schéma 2.
Illustration 2: Comparaison des puissances de calcul des CPU et des GPU
Cet écart est cependant à nuancer, la puissance de calcul des GPU donnée ici étant seulement en calcul flottant à 
simple précision et non en double comme pour les CPU. Le calcul en double précision n'est possible sur les GPU que depuis 
l'avènement du GT200 et reste pour le moment 8 fois plus lent que la puissance théorique en simple précision.
5L'écart de puissance théorique entre CPUs et GPUs s'explique par les natures intrinsèquement différentes des deux 
processeurs:
➢ Le CPU est optimisé pour exécuter le plus rapidement possible des séries de taches très variées. Il nécessite 
pour cela de structures de contrôles sophistiquées et doit disposer d'une mémoire cache pour réduire les coûts 
d'accès mémoire.
➢ Le GPU, de par sa nature graphique, est optimisé pour réaliser la même tache sur des milliers d'éléments 
différents( par exemple, un filtre sur une image va appliquer la même fonction à tous les pixels). Ainsi, ayant 
moins besoin d'unités de contrôles il peut comporter plus d'unités de calculs.
Illustration 3:Unités de calculs des CPU et GPU
a)  Modèle d'organisation des threads 
L'architecture des processeurs graphiques Nvidia Cuda est construite autour de multiprocesseurs (ou SM pour 
« Streaming Multiprocessor » en anglais ) multithreadés. Chaque SM est composé de 8 processeurs (ou SP pour « Streaming 
Processor » en anglais) qui réalisent les opérations de bases en simple précision flottante, auquel se rajoutent 2 unités 
spéciales (SFU) et 16kb de mémoire locale (appelée « shared memory »).
La puce T10 comporte en plus de cela une unité de calcul en double précision accolée à chaque SM. 
Illustration 5: Thread ProcessorIllustration 4: Streaming Multiprocessor de la T10
En pratique, ces SM sont regroupés par 3 sous formes de clusters (TPC pour « Thread Processing Cluster »). Ainsi, 
le T10 comporte 10 TPC, ce qui fait 10(TPC)*3(SM)*8(SP) = 240 coeurs de calculs.
Les   multiprocesseurs   ont   en   charge   la   création,   l'organisation   et   l'exécution   de   threads   sans   nécessité 
d'ordonnancement supplémentaire.
Pour cela, les SMs utilisent une architecture spécifique, appelée SIMT (pour « Single Instruction Multiple Thread ») 
par   Nvidia.   Chaque   thread   s'exécute   sur   un   processeur   indépendamment   des   autres   threads   avec   sa   propre   adresse 
d'exécution et son état de registre.
Plus précisément, les SMs partitionnent les blocs de threads de l'exécution (voir  Modèle de programmation) en 
groupe de 32 threads, appelés warps, qui seront ordonnancés par l'unité SIMT.
6A chaque temps d'instruction, l'unité de contrôle sélectionne un warp prêt à être exécuté, et réalise l'instruction 
suivante pour tous les threads du warp. Comme une seule instruction est réalisée à chaque fois, l'efficacité maximale est 
atteinte lorsque tous les threads du warp suivent le même chemin d'exécution. Dans le cas contraire, le temps d'exécution 
sera égale au nombre de branches divergentes prises par les threads.
Chaque   multiprocesseur   peut   gérer   jusqu'à   32   warps   actifs   simultanément,   ce   qui   fait   1024   threads   par 
multiprocesseur soit 30720 threads en exécution simultanée pour l'ensemble du GPU T10. Ce nombre important de threads 
permet aux unités de contrôle de masquer la latence des accès mémoire.
b)  Organisation hiérarchique de la mémoire  
La carte GPU embarque sa propre mémoire (4GB pour les cartes Tesla C1060), qui n'est pas adressable directement 
par le CPU et inversement pour la mémoire de la carte mère. Les données sont recopiées depuis le CPU et transitent via 
PCI express.
La mémoire vidéo lente rend les accès mémoire cruciaux en GPGPU il est important de bien en appréhender les 
contraintes.
Le schéma 6 permet d'avoir une vue globale de la hiérarchie de la mémoire.
La mémoire globale correspond à la mémoire centrale du GPU. Elle est accessible par tous les threads en lecture et 
en écriture mais aussi par le processeur hôte via l'api Cuda pour la recopie.
A cette mémoire se rajoute une zone de mémoire constante, qui est de la mémoire cache en lecture seule accessible 
par tous les coeurs de calcul. Il est également possible d'utiliser les unités de texturing du GPU comme une deuxième zone 
de mémoire cache en lecture seule. La différence principale étant que le cache de texture est optimisé pour des lectures 2D.
Localement, chaque processeur a à sa disposition quelques registres 32 bits, et une zone de mémoire cache, appelée 
shared memory, est commune à tous les coeurs d'un même multiprocesseur. 
Bien que la mémoire globale ait une grande latence, le contrôleur mémoire permet sous certaines conditions des 
accès mémoire groupés par demi warp (coalescing  en anglais). Ces conditions sont plus ou moins restrictives selon les 
générations de GPU (Compute Capability).
Pour le GPU T10 ( Compute Capability 1.3), les accès mémoire des 16 threads du demi warp sont regroupés en un 
seul accès dès que les mots machines accédés par les threads résident dans une zone mémoire contiguë de :
– 32 octets pour des mots machines de 8 bits
– 64 octets pour des mots machines de 16 bits
– 128 octets pour des mots machines de 32 ou 64 bits
Sinon le coût d'accès est égal au nombre de segments auxquels accède le demi wrap.
Comme la mémoire partagée est placée sur la puce, elle est beaucoup plus efficace que la mémoire globale. L'accès 
mémoire y est aussi rapide que l'accès aux registres tant qu'il n'y a pas de conflits de bancs.
Pour augmenter la bande passante mémoire, la mémoire partagée est divisée en 16 bancs qui sont des zones 
mémoires distinctes pouvant être accédées simultanément. Ce partage est fait de telle sorte que des mots machines successifs 
de 32­bits sont affiliés à des bancs successifs et que chaque banc ait une bande passante de 32 bits tous les 2 cycles.
On atteint donc le maximum de la bande passante possible lorsque tous les threads d'un demi warp accèdent à des 
bancs différents. Sinon, le temps d'accès est à priori le nombre maximal de threads accédant au même banc, sauf dans des 
cas de broadcasts supportés.
7Illustration 6: Hiérarchie de la mémoire dans le gpu
82. Modèle de programmation
Comme vu précédemment, le GPU est un coprocesseur extrêmement multithreadé avec sa propre mémoire. L'api C 
pour CUDA est une extension du C qui permet de s'adapter à cette architecture pour tirer profit de la puissance de calcul 
disponible.
Dans un programme, chaque phase de calcul massivement parallèle sur le GPU suit le schéma global suivant :
1. copie des données nécessaires au calcul dans la mémoire du GPU
2. exécution du même code N fois par N threads GPU en parallèle
3. recopie des résultats depuis la mémoire du GPU dans la mémoire centrale
a)  Kernels et threads 
Les kernels (noyaux en français) sont les fonctions qui spécifient le code qui sera exécuté par tous les threads en 
parallèle.
Un kernel est défini en utilisant le mot clé « _ _ global _ _  » :
//déclaration du kernel
__global__ void vecAdd(float* A, float* B, float* C){
//corps de la fonction
}
Un kernel est exécuté  sur une grille  (grid) de threads. L'appel d'un  kernel se fait toujours en spécifiant un 
environnement d'exécution entre chevrons qui décrit cette grille:
//invocation du kernel
vecAdd<<<grid, block>>>(A, B, C);
Une grille de threads est un ensemble bidimensionnel de blocs tridimensionnels de threads. Les paramètres de 
l'appel précédent sont les dimensions de la grille en nombre de blocs et celle des blocs en nombre de threads. Ces valeurs 
sont génériquement de type dim3, qui permet de décrire une dimension 3D: 
// configuration de l'environnement d'exécution
dim3 block(nbthreadx, nbthreads_y, nbthreads_z);
dim3 grid(nbblock_x, nbblock_y);
Un système d'index permet de repérer de façon unique chaque 
thread afin de spécifier sur quel lot de donnée il travaille.
Chaque thread est situé dans son bloc par des coordonnées 
tridimensionnelles via la variable ThreadIdx. De même chaque bloc est 
repéré dans la grille via un vecteur bidimensionnel BlockIdx. Enfin les 
variables d'environnement gridDim et blocDim permettent de connaître 
les dimensions respectivement de la grille et des blocs.
On repère ensuite intuitivement la colonne par .x, la ligne par 
.y, la profondeur par .z.
Illustration 7: Grille de threads
9On a donc deux niveaux de regroupement de threads:
– Un niveau « préemptif » : au niveau de la grille, l'ordre d'exécution des blocs est indéterminé et chaque bloc est exécuté 
complètement indépendamment des autres.
– Un niveau « coopératif » : tous les blocs ont le même nombre de threads (dans la limite de ce qu'accepte la carte, soit 
512 pour une Tesla C1060) et les threads d'un même bloc peuvent échanger des données via la mémoire partagée et se 
synchroniser à l'aide de la fonction _ _syncthreads().
b)  Utilisation de la mémoire 
L'allocation de mémoire sur le GPU s'effectue via la fonction cudaMalloc(), les transferts sont ensuite effectués à 
l'aide de la fonction cudaMemcpy().
float data[256],res[256];
int size = sizeof(data);
float* devPtr;
cudaMalloc((void**)&devPtr, size);
cudaMemcpy(devPtr, data, size, cudaMemcpyHostToDevice);
[⋅⋅⋅]//phase de calcul
cudaMemcpy(res, devPtr, size, cudaMemcpyDeviceToHost);
cudaFree(devPtr);
Pour   des   tableaux   multidimensionnels   il   est   préférable   d'utiliser   les   fonctions   cudaMallocPitch()   ou 
cudaMallocArray() qui permettent un meilleur alignement des données en mémoire, permettant ainsi plus d'accès groupés 
(coalesing, cf b) Organisation hiérarchique de la mémoire).
Toutes ces fonctions allouent de l'espace dans la mémoire globale mais, lors de l'exécution, chaque thread peut 
accéder aux différents niveaux de mémoire vus dans le chapitre précédent (cf b) Organisation hiérarchique de la mémoire). 
Dans le corps du kernel, les variables automatiques seront placées dans les registres et les déclarations utilisant le mot clé 
_ _shared_ _ seront dans la shared memory.
L'utilisation de la mémoire constante se fait en déclarant la variable au niveau global avec le mot clé _ _constant_ _ 
puis en copiant les données à l'aide de la fonction cudaMemcpyToSymbol(). Cependant, la zone de mémoire constante est 
assez petite (65Ko pour la Tesla C1060) donc son utilisation est restreinte à quelques constantes.
L'utilisation des unités de texturing est plus délicate. Elle nécessite de déclarer une texture au niveau global, de 
configurer une référence sur la texture, puis de lier la texture à des données en mémoire globale (voir exemple plus bas). Elle 
est recommandée lorsque les calculs voulus ne permettent aucun des paradigmes nécessaires au coalescing en mémoire 
globale.
//déclaration de la référence sur la texture
texture<float, 2, cudaReadModeElementType> M_texRef;
[...]
// création du CUDA array
cudaChannelFormatDesc chan_desc = cudaCreateChannelDesc<float>();
cudaArray *M_Array ;
//allocation
cudaMallocArray(&M_Array, &chan_desc,N,N);
//copie du la matrice M dans M_Array
cudaMemcpyToArray(M_Array,0,0,M,N*N*sizeof(float),cudaMemcpyHostToDevice);
//paramétrage de la texture
M_texRef.normalized = 0;
M_texRef.filterMode = cudaFilterModePoint;
M_texRef.addressMode[0] = cudaAddressModeClamp;
M_texRef.addressMode[1] = cudaAddressModeClamp;
//chaînage du tableau à la texture
cudaBindTextureToArray(M_texRef, M_Array,chan_desc);
[...]
//lecture de la texture 2D
r+=tex2D(M_texRef, (float)colonne_index, (float)ligne_index);
Exemple d'utilisation des unités de texturing
10