Nous travaillons sur le programme
prog00_hello.cpp.
Ce programme nous montrera que la mise en œuvre d'un GPU est très différente
de ce à quoi nous sommes habitués sur CPU.
Ce premier programme est complet ; il s'agit d'un exemple illustratif du
principe de fonctionnement d'un GPU (carte graphique) lorsqu'on veut
l'exploiter pour effectuer des calculs généralistes (
GP-GPU).
Vous n'aurez pas à le compléter mais il vous faudra l'étudier très
attentivement en suivant les explications.
Il existe différentes technologies pour le GP-GPU.
Le précurseur en la matière est l'environnement Nvidia CUDA : il s'agit
de la référence en termes de performances puisqu'il offre un contrôle
très fin du procédé mais il a l'inconvénient d'être spécifique au matériel
de la marque.
L'environnement OpenCL correspond quant à lui à un effort de normalisation
pour offrir des fonctionnalités similaires sur une plus grande variété de
matériel au prix d'un contrôle moins fin.
La technologie OpenGL, qui est à l'origine dédiée au rendu graphique 3D,
intègre depuis sa version 4.3 (en 2012) les
Compute-Shaders :
il s'agit de rendre accessibles les fonctionnalités GP-GPU dans le
contexte du rendu graphique afin de calculer sur le GPU les données qui
lui serviront lors du rendu.
Et il en existe d'autres :
https://en.wikipedia.org/wiki/General-purpose_computing_on_graphics_processing_units.
Quelle que soit la technologie choisie, elle repose toujours sur les mêmes
principes qui sont directement imposés par le fonctionnement de
l'architecture matérielle sous-jacente.
Dans le cadre de ces labos à l'ENIB, nous utiliserons l'API CUDA-driver
de Nvidia.
Il n'est pas question ici de passer en revue les détails de l'API retenue (qui
sont forcément différents pour une autre) c'est pourquoi le fichier
utilitaire
crsCuda.hpp les encapsule dans des objets de plus haut
niveau dont l'usage est suffisant pour la compréhension des principes
généraux.
Les fonctionnalités rendues accessibles sont très limitées mais suffisent
largement à la compréhension des éléments essentiels
Un usage plus classique pour débuter avec la technologie CUDA consiste
à utiliser le précompilateur nvcc fourni par cet environnement.
Il permet de rédiger avec le même formalisme du code qui concerne
le CPU et le GPU et se charge de le séparer automatiquement pour
le confier à chacun des deux compilateurs spécialisés tout en mettant
automatiquement en place les moyens de communication entre ces
deux moyens de calcul.
Si cela semble séduisant au premier abord, le confort ainsi apporté
est au détriment du contrôle fin de ce qui se déroule effectivement
dans le matériel, et dissimule les principes fondamentaux
sous-jacents.
Puisque cette expérimentation est à vocation pédagogique, nous nous
appuierons sur l'API CUDA-driver qui, elle au contraire, offre un
contrôle explicite de toutes les étapes du fonctionnement.
.
Commençons par le début du programme fourni.
Tout d'abord, le type utilitaire
crs::CudaPlatform sert à détecter
tout le matériel GPU compatible avec CUDA disponible sur la machine
informatique ; nous énumérons alors successivement chacun de ces
matériels en les désignant par la variable
device.
La fonction
to_string() rend visibles les propriétés qui ont été
détectées à l'initialisation de ce GPU.
Nous préparons également un paquet de données
hostData sur lesquelles
nous avons l'intention de faire travailler le GPU.
Il ne s'agit ici que de répéter une séquence d'entiers pour obtenir un
tableau assez grand ; ce n'est qu'un prétexte sans réelle utilité.
Vient ensuite le choix de la disposition à adopter pour effectuer le
traitement ; ceci est très lié à l'architecture matérielle.
Il faut imaginer un GPU comme constitué de plusieurs composants
multi-processeurs ; chacun d'eux est en effet capable d'exécuter
quasi-simultanément un grand nombre de traitements identiques.
Une vision simplifiée serait d'imaginer un grand nombre de
threads
exécutant tous la même instruction au même moment mais sur des données
distinctes pour chacun (modèle
SIMT, très différent du multi-CPU).
Un problème global doit alors être découpé en blocs ; chaque bloc est
confié à un tel composant multi-processeur.
Plusieurs composants travaillent simultanément sur des blocs distincts et
passent à un bloc suivant dès qu'ils en ont terminé un ; ainsi
l'ensemble des blocs constitutifs du problème sera globalement traité
par l'ensemble des composants multi-processeurs.
Plus le GPU dispose de tels composants, et plus ceux-ci peuvent exécuter
de traitements simultanés, plus massive sera la parallélisation du
traitement, ce qui doit conduire à un temps d'exécution très court par
rapport à une approche séquentielle.
Le choix de la taille des blocs et leur nombre est une question très
délicate qui, pour obtenir les performances optimales sur un matériel et un
problème donnés, nécessite de nombreuses mesures ; il n'existe pas de
réglage qui soit optimal pour tous les matériels et tous les problèmes ni
de règles automatiques pour les déterminer.
Nous nous contentons dans le cas présent de la fonction utilitaire
chooseLayout() qui détermine une valeur raisonnable pour ces deux
grandeurs.
Elle s'appuie sur les propriétés du GPU pour fournir des choix qui sont
connus pour donner des résultats généralement satisfaisants (mais pas
forcément optimaux) dans la plupart des cas.
Le code du traitement qui devra être exécuté sur le GPU est généré par
l'application (une simple chaîne de caractères) et compilé à la volée.
Nous voyons notamment que la fonction ainsi décrite porte le nom (
hello)
que nous avons choisi pour initialiser notre objet utilitaire
program.
Cette fonction est préfixée par
extern "C" __global__ ; c'est
obligatoire pour qu'elle soit retrouvée comme point d'entrée de ce
programme.
Son premier paramètre
data désigne un
buffer contenant des
entiers ; ce traitement pourra y écrire (il n'y a pas de
const
devant ce pointeur).
Un
buffer est un simple bloc de données situé dans la mémoire du GPU
(pas celle de la machine hôte !).
Deux autres paramètres entiers serviront à l'algorithme réalisé.
Il faut imaginer que chacun des
blockSize éléments des
blockCount
blocs est un
thread qui exécute cet algorithme.
Chacun d'eux est désigné de manière unique par la constante
globalId
dont la valeur sera comprise entre
0 et
gridSize-1 (
gridSize
vaut
blockSize*blockCount).
La détermination des constantes
gridSize et
globalId utilise ici
quelques variables prédéfinies du langage CUDA.
S'il y a plus de
gridSize éléments à traiter dans le problème il est
nécessaire que tous ces
threads fassent plusieurs itérations.
Puisque les
threads travaillent sur des indices contigus, ils doivent
progresser de
gridSize à chaque nouvelle itération.
Bien entendu, il faut veiller à ne pas dépasser le dernier élément du
problème (paramètre
count ici)
.
Le traitement retenu ici est trivial : nous nous contentons d'ajouter à
chaque élément du
buffer une valeur qui dépend de l'indice de calcul
et du paramètre
modulo.
Si tous les
threads effectuent cela à leurs indices respectifs, alors
tout le
buffer aura été traité
La syntaxe utilisée dans le code de ce traitement est bien entendu
spécifique à la technologie CUDA.
Elle serait différente pour une autre technologie mais les concepts
(buffers, paramètres, blocs, grille, identifiants...) seraient
similaires et elle reste généralement proche de celle du langage C.
.
Nous nous assurons enfin, avec la fonction utilitaire
assertSuccess(),
que le code en question est bien correct ; si ce n'est pas le cas, un
diagnostic sera donné et le programme sera terminé.
Dans le code, nous avons fait référence à un
buffer censé être présent
dans la mémoire du GPU.
Nous le créons alors en nous assurant que sa taille est suffisante
pour contenir toutes les données applicatives que nous avons produites
au début du programme.
Il est désormais envisageable de remplir le contenu de ce
buffer avec
les données en question.
Il s'agit de transférer, à travers le bus
PCI-express qui relie le GPU
à la carte-mère de la machine, le contenu du tableau en mémoire principale
(de la machine hôte) vers le
buffer dans la mémoire du GPU.
Une telle commande de transfert (ainsi que d'autres qui suivront) doit être
envoyée sur un flux de commandes qui est créé juste avant.
Maintenant que les données de notre problème sont censées avoir été recopiées
dans le
buffer du GPU, nous pouvons demander l'exécution du programme
sur GPU.
Avant cela, il nous faut néanmoins préparer un tableau de pointeurs sur des
données compatibles avec les paramètres attendus.
Le
buffer précédemment crée nous fournit un tel pointeur, et il suffit
de prendre l'adresse des deux entiers qui serviront à initialiser les deux
paramètres entiers suivants.
Cette commande d'exécution du programme sur GPU prend place dans le flux
de commandes, à la suite du transfert précédent.
L'exécution du programme sur le GPU est censée avoir eu pour effet de modifier
le contenu du
buffer que nous lui avons indiqué
; nous pouvons
dorénavant demander à en récupérer le contenu.
Il s'agit de transférer, à travers le bus
PCI-express qui relie le GPU
à la carte-mère de la machine, le contenu du
buffer dans la mémoire
du GPU vers le tableau en mémoire principale (de la machine hôte).
Cette commande de transfert prend place dans le flux de commandes, à la suite
de l'exécution précédente.
Désormais, les données produites par le GPU sont censées être utilisables par
la machine hôte.
Seulement le flux de commandes est asynchrone : ce n'est pas parce que nous
avons demandé la réalisation d'une commande que celle-ci a lieu
instantanément.
Il nous faut donc attendre jusqu'à ce que le flux de commandes nous indique
que la toute dernière commande de la séquence qui lui a été confiée est bien
terminée.
Nous nous contentons enfin d'afficher sous forme de texte les données
reçues depuis le GPU.
Pour tester ce programme, il suffit de le lancer sans plus d'arguments.
Vous devriez voir apparaître dans le terminal les propriétés du GPU et les
données produites par l'exécution du programme que nous lui avons confié.
Assurez vous d'avoir bien compris le rôle de toutes les étapes importantes du
procédé qui ont été décrites ici :
- fabrication d'un programme pour le GPU, organisé en blocs et exécuté en
parallèle par de multiples threads désignés par un identifiant
entier,
- création de buffers sur le GPU,
- transferts de données tableau-hôte → buffer-GPU,
- exécution du programme sur GPU pour exploiter les buffers et en
considérant des paramètres,
- transferts de données buffer-GPU → tableau-hôte.
Ce sont en effet les éléments incontournables qui sont à la base de toute
application qui exploite un GPU.
Documentation utile :