Chapitre 2 : Modèles de programmation GPU
Les unités de traitement graphique (GPU) ont évolué des accélérateurs graphiques à fonction fixe vers des moteurs de calcul hautement parallèles et programmables capables d'accélérer une large gamme d'applications. Pour permettre aux programmeurs d'exploiter efficacement le parallélisme massif des GPU, plusieurs modèles de programmation parallèle et API ont été développés, comme NVIDIA CUDA, OpenCL et DirectCompute. Ces modèles de programmation fournissent des abstractions qui permettent aux programmeurs d'exprimer le parallélisme dans leurs applications tout en masquant les détails de bas niveau du matériel GPU.
Dans ce chapitre, nous explorerons les concepts et principes clés derrière les modèles de programmation parallèle pour les GPU, en nous concentrant sur le modèle d'exécution, les architectures d'ensemble d'instructions (ISA) des GPU, les ISA des GPU NVIDIA et l'ISA Graphics Core Next (GCN) d'AMD. Nous fournirons également des exemples pour illustrer comment ces concepts sont appliqués dans la pratique.
Modèle d'exécution
Le modèle d'exécution des modèles de programmation GPU modernes est basé sur le concept de noyaux, qui sont des fonctions exécutées en parallèle par un grand nombre de threads sur le GPU. Lors du lancement d'un noyau, le programmeur spécifie le nombre de threads à créer et la manière dont ils sont organisés dans une hiérarchie de grilles, de blocs (ou tableaux de threads coopératifs - CTA) et de threads individuels.
- Une grille représente l'espace du problème dans son ensemble et se compose d'un ou plusieurs blocs.
- Un bloc est un groupe de threads qui peuvent coopérer et se synchroniser entre eux via la mémoire partagée et les barrières. Les threads d'un même bloc sont exécutés sur le même cœur GPU (appelé multiprocesseur de flux ou unité de calcul).
- Chaque thread a un ID unique dans son bloc et sa grille, ce qui peut être utilisé pour calculer les adresses mémoire et prendre des décisions de flux de contrôle.
Cette organisation hiérarchique permet aux programmeurs d'exprimer à la fois le parallélisme de données (où la même opération est appliquée à plusieurs éléments de données) et le parallélisme de tâches (où différentes tâches sont exécutées en parallèle).
FigurVoici la traduction française du fichier markdown, avec les commentaires traduits mais le code non traduit :
e 2.1 illustre la hiérarchie des threads dans le modèle d'exécution GPU.
Grid
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Block |
| | | |
Thread Thread ...Figure 2.1 : Hiérarchie des threads dans le modèle d'exécution GPU.
Exécution SIMT
Les modèles de programmation GPU comme CUDA et OpenCL suivent un modèle d'exécution à Instruction Unique, Multiples Threads (SIMT). Dans le modèle SIMT, les threads sont exécutés en groupes appelés warps (terminologie NVIDIA) ou wavefronts (terminologie AMD). Tous les threads au sein d'un warp exécutent la même instruction en même temps, mais chaque thread opère sur des données différentes.
Cependant, contrairement au modèle traditionnel à Instruction Unique, Données Multiples (SIMD), où tous les éléments de traitement s'exécutent en lock-step, SIMT permet aux threads d'avoir des chemins d'exécution indépendants et de diverger aux instructions de branchement. Lorsqu'un warp rencontre une instruction de branchement, le matériel GPU évalue la condition de branchement pour chaque thread du warp. Si tous les threads empruntent le même chemin (convergent), le warp continue l'exécution normalement. Si certains threads empruntent des chemins différents (divergents), le warp est divisé en deux sous-warps ou plus, chacun suivant un chemin différent. Le matériel GPU sérialise l'exécution des chemins divergents, masquant les threads inactifs dans chaque sous-warp. Lorsque tous les chemins sont terminés, les sous-warps reconvergent et continuent l'exécution en lock-step.
La figure 2.2 illustre l'exécution SIMT avec un flux de contrôle divergent.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Branch |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
\
\
ReconvergenceFigure 2.2 : Exécution SIMT avec un flux de contrôle divergent.
Ce mécanisme de gestion de la divergence permet à SIMT de prendre en charge un flux de contrôle plus flexible tVoici la traduction française du fichier Markdown, avec les commentaires du code traduits, mais le code lui-même non modifié :
Hiérarchie de la mémoire
Les GPU ont une hiérarchie de mémoire complexe pour prendre en charge les exigences de bande passante élevée et de faible latence des charges de travail parallèles. La hiérarchie de la mémoire se compose généralement de :
- Mémoire globale : l'espace mémoire le plus grand mais le plus lent, accessible par tous les threads dans un noyau. La mémoire globale est généralement mise en œuvre à l'aide de mémoire GDDR ou HBM à bande passante élevée.
- Mémoire partagée : un espace mémoire rapide et sur puce partagé par tous les threads d'un bloc. La mémoire partagée est utilisée pour la communication inter-threads et le partage de données au sein d'un bloc.
- Mémoire constante : un espace mémoire en lecture seule utilisé pour diffuser des données en lecture seule à tous les threads.
- Mémoire de texture : un espace mémoire en lecture seule optimisé pour la localité spatiale et accessible via les caches de texture. La mémoire de texture est plus couramment utilisée dans les charges de travail graphiques.
- Mémoire locale : un espace mémoire privé pour chaque thread, utilisé pour le débordement des registres et les grandes structures de données. La mémoire locale est généralement mappée à la mémoire globale.
Une utilisation efficace de la hiérarchie de la mémoire est cruciale pour atteindre de hautes performances sur les GPU. Les programmeurs doivent viser à maximiser l'utilisation de la mémoire partagée et à minimiser les accès à la mémoire globale pour réduire la latence de la mémoire et les goulots d'étranglement de la bande passante.
La figure 2.3 illustre la hiérarchie de la mémoire GPU.
____________
| |
| Global |
| Memory |
____________
|
____________
| |
| Constant |
| Memory |
____________
|
____________
| |
| Texture |
| Memory |
____________
|
|
____________
| |
| Shared |
| Memory |
____________
|
____________
| |
| Local |
| Memory |
____________FigVoici la traduction française du fichier markdown, avec les commentaires traduits pour le code, mais sans traduction du code lui-même :
Architectures d'ensembles d'instructions GPU
Les architectures d'ensembles d'instructions GPU (ISA) définissent l'interface de bas niveau entre le logiciel et le matériel. Elles spécifient les instructions, les registres et les modes d'adressage mémoire pris en charge par le GPU. Comprendre les ISA des GPU est essentiel pour développer du code GPU efficace et optimiser les performances.
Dans cette section, nous explorerons les ISA de deux principaux fournisseurs de GPU : NVIDIA et AMD. Nous nous concentrerons sur les ISA PTX (Parallel Thread Execution) et SASS d'NVIDIA, ainsi que sur l'ISA Graphics Core Next (GCN) d'AMD.
ISA des GPU NVIDIA
Les GPU NVIDIA prennent en charge deux niveaux d'ISA : PTX (Parallel Thread Execution) et SASS (Streaming ASSembler). PTX est une ISA virtuelle qui fournit une cible stable pour les compilateurs CUDA, tandis que SASS est l'ISA native des GPU NVIDIA.
PTX (Parallel Thread Execution)
PTX est une ISA virtuelle de bas niveau conçue pour les GPU NVIDIA. Elle est similaire à LLVM IR ou Java bytecode dans le sens où elle fournit une cible stable et indépendante de l'architecture pour les compilateurs. Les programmes CUDA sont généralement compilés en code PTX, qui est ensuite traduit en instructions SASS natives par le pilote GPU NVIDIA.
PTX prend en charge une large gamme d'instructions arithmétiques, mémoire et de contrôle de flux. Elle dispose d'un nombre illimité de registres virtuels et prend en charge la prédication, ce qui permet une mise en œuvre efficace du contrôle de flux. PTX fournit également des instructions spéciales pour la synchronisation des threads, les opérations atomiques et l'échantillonnage de texture.
Voici un exemple de code PTX pour un noyau d'addition de vecteurs simple :
.version 7.0
.target sm_70
.address_size 64
.visible .entry vecAdd(
.param .u64 vecAdd_param_0,
.param .u64 vecAdd_param_1,
.param .u64 vecAdd_param_2,
.param .u32 vecAdd_param_3
)
{
.reg .b32 %r<4>;
.reg .b64 %rd<6>;
// Charger les paramètres
ld.param.u64 %rd1, [vecAdd_param_0];
ld.param.u64 %rd2, [vecAdd_param_1];
ld.param.u64 %rd3, [vecAdd_param_2];
ld.param.u32 %r1, [vecAdd_param_3];
// Convertir les adresses en adresses globales
cvta.to.global.u64 %rd4, %rd1;
cvta
```Voici la traduction française du fichier Markdown avec les commentaires traduits, mais sans traduire le code :
// .to.global.u64 %rd5, %rd2; // mov.u32 %r2, %tid.x; // mul.wide.u32 %rd6, %r2, 4; // add.s64 %rd7, %rd4, %rd6; // add.s64 %rd8, %rd5, %rd6;
// ld.global.f32 %f1, [%rd7]; // ld.global.f32 %f2, [%rd8]; // add.f32 %f3, %f1, %f2;
// cvta.to.global.u64 %rd9, %rd3; // add.s64 %rd10, %rd9, %rd6; // st.global.f32 [%rd10], %f3;
// ret; }
Ce code PTX définit une fonction de noyau `vecAdd` qui prend quatre paramètres : des pointeurs vers les vecteurs d'entrée et de sortie, et la taille des vecteurs. Le noyau calcule l'ID de thread global, charge les éléments correspondants à partir des vecteurs d'entrée, effectue l'addition et stocke le résultat dans le vecteur de sortie.
#### SASS (Streaming ASSembler)
SASS est l'ISA native des cartes graphiques NVIDIA. C'est une ISA de bas niveau, spécifique à la machine, qui correspond directement au matériel GPU. Les instructions SASS sont générées par le pilote GPU NVIDIA à partir du code PTX et ne sont généralement pas visibles pour les programmeurs.
Les instructions SASS sont codées dans un format compact pour réduire la bande passante mémoire et l'empreinte du cache d'instructions. Elles prennent en charge une large gamme de types d'opérandes, notamment des registres, des valeurs immédiates et divers modes d'adressage pour l'accès à la mémoire.
Voici un exemple de code SASS pour le noyau d'addition de vecteurs :
```sass
code_version_number 90
@P0 LDG.E R2, [R8];
@P0 LDG.E R0, [R10];
@P0 FADD R0, R2, R0;
@P0 STG.E [R12], R0;
EXIT;Ce code SASS correspond au code PTX présenté précédemment. Il charge les éléments des vecteurs d'entrée depuis la mémoire globale (LDG.E), effectue l'addition (FADD), stocke le résultat dans la mémoire globale (STG.E) et sort du noyau.
AMD Graphics Core Next ISA
Les cartes graphiques AMD utilisent l'architecture et l'ISA Graphics Core Next (GCN). GCN est une ISA RISC conçue pour prendre en charge à la fois les charges de travail graphiques et de calcul. Elle est conçue pour offrir de hautes performances, une grande extensibilité et une efficacité énergétique élevée.
GCN introduit plusieurs fonctionnalités clés, comme :
- UneVoici la traduction française du fichier Markdown avec la traduction des commentaires du code GCN, mais sans traduire le code lui-même :
alar ALU pour une exécution efficace des opérations scalaires et du contrôle de flux.
- Une ALU vectorielle pour l'exécution parallèle d'opérations parallèles sur les données.
- Un système de mémoire à bande passante élevée avec prise en charge des opérations atomiques et un accès à faible latence à la mémoire partagée.
- Un mode d'adressage flexible pour les opérations mémoire, prenant en charge l'adressage base+décalage et scalaire+vectoriel.
Voici un exemple de code ISA GCN pour un noyau d'addition de vecteurs :
.text
.globl vecAdd
.p2align 2
.type vecAdd,@function
vecAdd:
.set DPTR, 0
# Charger les arguments du noyau à partir de la mémoire
s_load_dwordx4 s[0:3], s[4:5], 0x0
s_load_dword s4, s[4:5], 0x10
s_waitcnt lgkmcnt(0)
# Calculer l'adresse de l'élément du vecteur
v_lshlrev_b32 v0, 2, v0
v_add_u32 v1, vcc, s1, v0
v_mov_b32 v3, s3
v_addc_u32 v2, vcc, s2, v3, vcc
flat_load_dword v1, v[1:2]
# Calculer l'adresse de l'élément du vecteur
v_add_u32 v3, vcc, s0, v0
v_mov_b32 v5, s3
v_addc_u32 v4, vcc, s2, v5, vcc
flat_load_dword v0, v[3:4]
# Effectuer l'addition des vecteurs
v_add_f32 v0, v0, v1
flat_store_dword v[3:4], v0
s_endpgmMappage des algorithmes aux architectures GPU
Le mappage efficace des algorithmes à l'architecture GPU est essentiel pour obtenir de hautes performances. Les principaux éléments à prendre en compte sont les suivants :
Exposer suffisamment de parallélisme
L'algorithme doit être décomposé en de nombreux fils d'exécution de grain fin qui peuvent s'exécuter en parallèle pour utiliser pleinement les capacités de traitement parallèle du GPU. Cela implique souvent d'identifier les parties de l'algorithme parallèles sur les données qui peuvent être exécutées indépendamment sur différents éléments de données.
Minimiser la divergence de branches
La divergence du flux de contrôle au sein d'un warp/wavefront peut entraîner une sérialisation et une efficacité SIMD réduite. Les algorithmes doivent être structurés pour minimiser la divergence des branches autant que possible. Cela peut être réalisé en réduisant l'utilisation de flux de contrôle dépendants des données.Voici la traduction française du fichier markdown :
Exploiter la hiérarchie de la mémoire
Accéder à la mémoire globale est coûteux. Les algorithmes doivent maximiser l'utilisation de la mémoire partagée et des registres pour réduire les accès à la mémoire globale. Les données doivent également être organisées en mémoire pour permettre des accès mémoire groupés, où les threads d'un warp accèdent à des emplacements de mémoire contigus. Une utilisation efficace de la hiérarchie de la mémoire peut réduire considérablement la latence de la mémoire et les goulots d'étranglement de la bande passante.
Équilibrer le calcul et les accès à la mémoire
Les algorithmes doivent avoir un rapport élevé d'opérations arithmétiques par rapport aux opérations de mémoire pour masquer efficacement la latence de la mémoire et atteindre un débit de calcul élevé. Cela peut être obtenu en maximisant la réutilisation des données, en préchargeant les données et en chevauchant les calculs avec les accès à la mémoire.
Minimiser les transferts de données hôte-périphérique
Le transfert de données entre la mémoire de l'hôte (CPU) et la mémoire du périphérique (GPU) est lent. Les algorithmes doivent minimiser ces transferts en effectuant autant de calculs que possible sur le GPU. Les données doivent être transférées vers le GPU par lots importants et conservées sur le périphérique aussi longtemps que nécessaire pour amortir les coûts de transfert.
Plusieurs modèles de conception d'algorithmes parallèles sont couramment utilisés lors du développement de noyaux GPU :
-
Map : Chaque thread effectue la même opération sur un élément de données différent, permettant un traitement parallèle simple de grands ensembles de données.
-
Reduce : La réduction parallèle est utilisée pour calculer efficacement une seule valeur (par exemple, la somme, le maximum) à partir d'un grand ensemble de données d'entrée. Les threads effectuent des réductions locales, qui sont ensuite combinées pour produire le résultat final.
-
Scan : Également appelé préfixe, le scan sert à calculer la somme cumulée des éléments d'un tableau. Les algorithmes de scan parallèles efficaces sont des éléments de base essentiels pour de nombreuses applications accélérées par GPU.
-
Stencil : Chaque thread calcule une valeur en fonction des éléments de données voisins. Les calculs de stencil sont courants dans les simulations scientifiques et le traitement d'images.Voici la traduction française du fichier markdown :
-
Gather/Scatter : Les threads lisent (gather) ou écrivent (scatter) à des emplacements arbitraires dans la mémoire globale. Une disposition et des motifs d'accès aux données soignés sont nécessaires pour être efficace.
La figure 3.20 illustre un exemple du modèle de mappage, où chaque thread applique une fonction (par exemple, la racine carrée) à un élément différent du tableau d'entrée.
Tableau d'entrée :
| | | | | | | |
v v v v v v v v
______________________________
Threads : | | | | | | | |
|____|____|____|____|____|____|____|
| | | | | | |
v v v v v v v
Tableau de sortie : Figure 3.20 : Exemple du modèle de mappage dans la programmation GPU.
Conclusion
Les modèles de programmation GPU comme CUDA et OpenCL exposent les capacités de traitement parallèle des GPU modernes aux développeurs, leur permettant d'accélérer une grande variété d'applications. Ces modèles de programmation fournissent des abstractions qui permettent de mapper efficacement les charges de travail parallèles à grain fin sur le matériel GPU.
Comprendre le modèle d'exécution, la hiérarchie de la mémoire et les primitives de synchronisation fournis par ces modèles de programmation est essentiel pour écrire du code GPU haute performance. Les développeurs doivent prendre en compte des facteurs tels que l'organisation des threads, la divergence des branches, les motifs d'accès à la mémoire et la conception des algorithmes pour exploiter pleinement la puissance de calcul des GPU.
À mesure que les architectures GPU continuent d'évoluer, les modèles de programmation et les outils doivent également progresser pour permettre aux développeurs d'utiliser efficacement les nouvelles fonctionnalités et capacités du matériel. La recherche continue dans des domaines tels que la conception de langages de programmation, l'optimisation des compilateurs et l'auto-ajustement sera cruciale pour améliorer la productivité des programmeurs et la portabilité des performances à l'ère du calcul hétérogène.