Chapitre 3 : Modèles de programmation parallèle dans la conception de GPU
Les unités de traitement graphique (GPU) ont évolué à partir d'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, tels que 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 SIMT (Single Instruction, Multiple Thread), le modèle de programmation CUDA et les API, ainsi que le cadre OpenCL. Nous aborderons également les techniques permettant de mapper les algorithmes sur les architectures GPU pour atteindre des performances et une efficacité élevées.
Modèle d'exécution SIMT (Single Instruction, Multiple Thread)
Le modèle d'exécution SIMT est le paradigme fondamental utilisé par les GPU modernes pour réaliser un parallélisme massif. Dans le modèle SIMT, un grand nombre de threads exécutent le même programme (appelé noyau) en parallèle, mais chaque thread a son propre compteur de programme et peut emprunter différents chemins d'exécution en fonction de son identifiant de thread et des données sur lesquelles il opère.
Noyaux et hiérarchie des threads
Un noyau GPU est une fonction qui est exécutée en parallèle par un grand nombre de threads. Lors du lancement d'un noyau, le programmeur spécifie le nombre de threads à créer et comment 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 au sein d'un bloc sont exécutés sur le même cœur GPU (appelé multiprocesseur de flux).Voici la traduction française du fichier Markdown, avec les commentaires traduits, mais le code laissé inchangé :
(ou unité de calcul).
- Chaque thread a un identifiant unique au sein de son bloc et de sa grille, 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).
La figure 3.1 illustre la hiérarchie des threads dans le modèle d'exécution SIMT.
Grille
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Bloc |
| | | |
Thread Thread ...Figure 3.1 : Hiérarchie des threads dans le modèle d'exécution SIMT.
Exécution SIMT
Dans le modèle d'exécution SIMT, chaque thread exécute la même instruction mais opère sur des données différentes. Cependant, contrairement au SIMD (Single Instruction, Multiple Data) où tous les éléments de traitement s'exécutent de manière synchrone, le SIMT permet aux threads d'avoir des chemins d'exécution indépendants et de diverger aux instructions de branchement.
Lorsqu'un warp (un groupe de 32 threads sur les GPU NVIDIA ou 64 threads sur les GPU AMD) 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 (convergé), le warp continue l'exécution normalement. Cependant, si certains threads empruntent des chemins différents (divergés), le warp est divisé en deux ou plusieurs sous-warps, 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 se reconvergent et continuent l'exécution de manière synchrone.
La figure 3.2 illustre l'exécution SIMT avec un flux de contrôle divergent.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Branchement |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
```Reconvergence
Figure 3.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 que SIMD, mais cela se fait au détriment d'une efficacité SIMD réduite lorsque la divergence se produit. Les programmeurs doivent s'efforcer de minimiser la divergence au sein d'un warp pour obtenir des performances optimales.
Hiérarchie de la mémoire
Les GPU ont une hiérarchie de mémoire complexe pour répondre aux exigences élevées de bande passante 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 haute bande passante.
- 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-thread 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 des 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 de registre et les structures de données de grande taille. La mémoire locale est généralement mappée sur la mémoire globale.
L'utilisation efficace de la hiérarchie de la mémoire est cruciale pour atteindre des performances élevées 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 afin de réduire les goulots d'étranglement de latence et de bande passante mémoire.
La figure 3.3 illustre la hiérarchie de la mémoire GPU.
| Mémoire |
| Partagée |
____________
|
____________
| |
| Mémoire |
| Locale |
____________Figure 3.3 : Hiérarchie de la mémoire GPU.
Modèle de programmation CUDA et API
CUDA (Compute Unified Device Architecture) est une plateforme de calcul parallèle et un modèle de programmation développés par NVIDIA pour le calcul généraliste sur GPU. CUDA fournit un ensemble d'extensions aux langages de programmation standard, tels que C, C++ et Fortran, qui permettent aux programmeurs d'exprimer le parallélisme et d'exploiter la puissance de calcul des GPU NVIDIA.
Modèle de programmation CUDA
Le modèle de programmation CUDA repose sur le concept de noyaux, qui sont des fonctions exécutées en parallèle par un grand nombre de threads sur le GPU. Le programmeur spécifie le nombre de threads à lancer et leur organisation en une grille de blocs de threads.
CUDA introduit plusieurs abstractions clés pour faciliter la programmation parallèle :
- Thread : L'unité de base d'exécution dans CUDA. Chaque thread a son propre compteur de programme, ses propres registres et sa propre mémoire locale.
- Bloc : Un groupe de threads qui peuvent coopérer et se synchroniser entre eux. Les threads au sein d'un bloc sont exécutés sur le même multiprocesseur de flux et peuvent communiquer via la mémoire partagée.
- Grille : Une collection de blocs de threads qui exécutent le même noyau. La grille représente l'espace du problème complet et peut être unidimensionnelle, bidimensionnelle ou tridimensionnelle.
CUDA fournit également des variables intégrées (par exemple, threadIdx, blockIdx, blockDim, gridDim) qui permettent aux threads de s'identifier et de calculer des adresses mémoire en fonction de leur position dans la hiérarchie des threads.
La figure 3.4 illustre le modèle de programmation CUDA.
Grille
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Bloc |
| | | |
Thread Thread ...Figure 3.4 : Modèle de programmation CUDA.
Hiérarchie de la mémoire CUDAVoici la traduction française du fichier Markdown "archy" avec les commentaires traduits, mais le code reste inchangé.
CUDA expose la hiérarchie de la mémoire GPU au programmeur, permettant un contrôle explicite sur le placement et le déplacement des données. Les principaux espaces mémoire dans CUDA sont :
- Mémoire globale : Accessible à tous les threads dans un noyau et persistent entre les lancements de noyaux. La mémoire globale a la latence la plus élevée et est généralement utilisée pour les grandes structures de données.
- Mémoire partagée : Une mémoire rapide, intégrée sur puce, partagée 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. La mémoire constante est mise en cache et offre un accès à faible latence.
- 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.
Les programmeurs peuvent allouer et transférer des données entre la mémoire hôte (CPU) et la mémoire de l'appareil (GPU) en utilisant les API runtime CUDA, telles que cudaMalloc, cudaMemcpy et cudaFree.
La figure 3.5 illustre la hiérarchie de la mémoire CUDA.
____________
| |
| Global |
| Memory |
____________
|
____________
| |
| Constant |
| Memory |
____________
|
____________
| |
| Texture |
| Memory |
____________
|
|
____________
| |
| Shared |
| Memory |
____________
|
____________
| |
| Local |
| Memory |
____________Figure 3.5 : Hiérarchie de la mémoire CUDA.
Synchronisation et coordination CUDA
CUDA fournit des primitives de synchronisation et de coordination pour permettre la coopération et la communication entre les threads :
- Synchronisation par barrière : La fonction __syncthreads()Voici la traduction française de ce fichier Markdown :
La fonction s() agit comme une barrière qui s'assure que tous les fils d'exécution d'un bloc ont atteint le même point avant de poursuivre.
- Opérations atomiques : CUDA prend en charge les opérations atomiques (par exemple, atomicAdd, atomicExch) qui permettent aux fils d'exécution d'effectuer des opérations de lecture-modification-écriture sur la mémoire partagée ou globale sans interférence d'autres fils d'exécution.
- Primitives de niveau warp : CUDA fournit des intrinsèques de niveau warp (par exemple, __shfl, __ballot) qui permettent une communication et une synchronisation efficaces au sein d'un warp.
L'utilisation appropriée des primitives de synchronisation et de coordination est essentielle pour écrire des programmes parallèles corrects et efficaces en CUDA.
L'exemple 3.1 montre un noyau CUDA simple qui effectue l'addition de vecteurs.
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
c[i] = a[i] + b[i];
}
}
int main() {
int *a, *b, *c;
int n = 1024;
// Allouer la mémoire sur l'hôte
a = (int*)malloc(n * sizeof(int));
b = (int*)malloc(n * sizeof(int));
c = (int*)malloc(n * sizeof(int));
// Initialiser les vecteurs d'entrée
for (int i = 0; i < n; i++) {
a[i] = i;
b[i] = i * 2;
}
// Allouer la mémoire sur le périphérique
int *d_a, *d_b, *d_c;
cudaMalloc(&d_a, n * sizeof(int));
cudaMalloc(&d_b, n * sizeof(int));
cudaMalloc(&d_c, n * sizeof(int));
// Copier les vecteurs d'entrée de l'hôte vers le périphérique
cudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice);
// Lancer le noyau
int blockSize = 256;
int numBlocks = (n + blockSize - 1) / blockSize;
vectorAdd<<<numBlocks,blockSize>>>(d_a, d_b, d_c, n);
// Copier le vecteur de résultat du périphérique vers l'hôte
cudaMemcpy(c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);
// Libérer la mémoire du périphérique
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
// Libérer la mémoire de l'hôte
free(a);
free(b);
free(c);
returVoici la traduction française du fichier Markdown, avec les commentaires traduits mais le code non traduit :
```c
int main() {
// Initialisation des variables
int numBlocks = 16;
int blockSize = 256;
int size = numBlocks * blockSize * sizeof(float);
// Allocation de mémoire sur l'hôte
float *a = (float*)malloc(size);
float *b = (float*)malloc(size);
float *c = (float*)malloc(size);
// Initialisation des vecteurs d'entrée
for (int i = 0; i < numBlocks * blockSize; i++) {
a[i] = 1.0f;
b[i] = 2.0f;
}
// Allocation de mémoire sur le périphérique
float *d_a, *d_b, *d_c;
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
// Copie des données de l'hôte vers le périphérique
cudaMemcpy(d_a, a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, size, cudaMemcpyHostToDevice);
// Lancement du noyau CUDA
vectorAdd<<<numBlocks, blockSize>>>(d_a, d_b, d_c);
// Copie des résultats du périphérique vers l'hôte
cudaMemcpy(c, d_c, size, cudaMemcpyDeviceToHost);
// Libération de la mémoire
free(a);
free(b);
free(c);
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
return 0;
}Ce code CUDA lance le noyau vectorAdd avec numBlocks blocs et blockSize threads par bloc. Le noyau effectue une addition élément par élément des vecteurs d'entrée a et b et stocke le résultat dans le vecteur c. La syntaxe <<<...>>> est utilisée pour spécifier les dimensions de la grille et des blocs lors du lancement d'un noyau.
Flux et événements CUDA
Les flux et événements CUDA fournissent un mécanisme pour l'exécution concurrente et la synchronisation des noyaux et des opérations de mémoire :
- Flux : Une séquence d'opérations (lancements de noyaux, copies de mémoire) qui s'exécutent dans l'ordre. Différents flux peuvent s'exécuter de manière concurrente, permettant le chevauchement des calculs et des transferts de mémoire.
- Événements : Des marqueurs qui peuvent être insérés dans un flux pour enregistrer l'achèvement d'opérations spécifiques. Les événements peuvent être utilisés à des fins de synchronisation et de chronométrage.
Les flux et les événements permettent aux programmeurs d'optimiser les performances de leurs applications CUDA en chevauchant les calculs et les transferts de mémoire, et en exploitant pleinement les capacités du matériel GPU.
L'exemple 3.2 montre l'utilisation des flux CUDA pour chevaucher l'exécution des noyaux et les transferts de mémoire.
// Création de deux flux
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// Copie asynchrone des données d'entrée vers le périphérique
cudaMemcpyAsync(d_a, a, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, b, size, cudaMemcpyHostToDevice, stream2);
// Lancement des noyaux dans différents flux
kernelA<<<blocks, threads, 0, stream1>>>(d_a);
kernelB<<<blocks, threads, 0, stream2>>>(d_b);
// Copie asynchrone des résultats de retour vers l'hôte
cudaMemcpyAsync(a, d_a, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(b, d_b, size, cudaMemcpyDeviceToHost, stream2);
// Synchronisation des flux
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);Dans cet exemple, deux flux CUDA sont créés. Les données d'entrée sont copiées de manière asynchrone vers le périphérique en utilisant chaque flux. Ensuite, les noyaux sont lancés dans les différents flux, permettant le chevauchement des calculs et des transferts de mémoire.Voici la traduction française du fichier Markdown, avec les commentaires du code traduits en français :
Cadre OpenCL
OpenCL (Open Computing Language) est une norme ouverte et sans redevance pour la programmation parallèle sur des plateformes hétérogènes, y compris les processeurs, les cartes graphiques, les FPGA et d'autres accélérateurs. OpenCL fournit un modèle de programmation unifié et un ensemble d'API qui permettent aux développeurs d'écrire du code parallèle portable et efficace.
Modèle de programmation OpenCL
Le modèle de programmation OpenCL est similaire à CUDA, avec quelques différences clés dans la terminologie et les abstractions :
- Kernel : Une fonction exécutée en parallèle par un grand nombre d'éléments de travail (threads) sur un dispositif OpenCL.
- Élément de travail : L'unité de base d'exécution dans OpenCL, analogue à un thread dans CUDA.
- Groupe de travail : Un ensemble d'éléments de travail qui peuvent se synchroniser et partager des données via la mémoire locale. Les groupes de travail sont analogues aux blocs de threads dans CUDA.
- Espace N-dimensionnel (NDRange) : Définit l'espace d'indexation et l'organisation des éléments de travail pour l'exécution d'un kernel. Il peut être à une, deux ou trois dimensions.
OpenCL définit également un modèle de mémoire hiérarchique similaire à CUDA :
- Mémoire globale : Accessible par tous les éléments de travail dans tous les groupes de travail, analogue à la mémoire globale dans CUDA.
- Mémoire locale : Partagée par tous les éléments de travail d'un groupe de travail, analogue à la mémoire partagée dans CUDA.
- Mémoire privée : Privée à un seul élément de travail, analogue aux registres dans CUDA.
- Mémoire constante : Mémoire en lecture seule accessible par tous les éléments de travail.
Les kernels OpenCL sont compilés à l'exécution par le runtime OpenCL. Le programme hôte peut interroger les périphériques OpenCL disponibles, sélectionner un périphérique approprié, créer un contexte et compiler le kernel pour ce périphérique spécifique. Cela permet aux applications OpenCL d'être hautement portables sur différentes plateformes matérielles.
L'exemple 3.3 montre un kernel OpenCL qui effectue une addition de vecteurs, similaire à l'exemple CUDA de l'exemple 3.1.
__kernel void vectorAdd(__global const int *a, __global int *b, __global int *c) {
// Récupérer l'identifiant unique de l'élément de travail
int i = get_global_id(0);
// Effectuer l'opération d'addition de vecteurs
c[i] = a[i] + b[i];
}
```Voici la traduction en français de ce fichier Markdown, avec les commentaires traduits mais le code source laissé inchangé :
```c
__kernel void bal const int *b, __global int *c, int n) {
int i = get_global_id(0);
if (i < n) {
c[i] = a[i] + b[i];
}
}Le mot-clé __kernel définit une fonction de noyau OpenCL. Le mot-clé __global indique qu'un pointeur pointe vers la mémoire globale. La fonction get_global_id renvoie l'index global de l'unité de travail actuelle, qui est utilisé pour calculer les adresses mémoire des vecteurs d'entrée et de sortie.
Mappage des algorithmes aux architectures GPU
Le mappage efficace des algorithmes à l'architecture GPU est crucial pour atteindre des performances élevées. Les principales considérations sont les suivantes :
-
Exposer suffisamment de parallélisme : L'algorithme doit être décomposé en de nombreux threads de grain fin qui peuvent s'exécuter simultanément pour utiliser pleinement les capacités de traitement parallèle du GPU.
-
Minimiser la divergence des 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.
-
Exploiter la hiérarchie de la mémoire : L'accès à 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 coalescents.
-
Équilibrer les calculs et les accès mémoire : Les algorithmes doivent avoir un ratio élevé d'opérations arithmétiques par rapport aux opérations mémoire pour masquer efficacement la latence mémoire et atteindre un débit de calcul élevé.
-
Minimiser les transferts de données hôte-périphérique : Le transfert de données entre la mémoire hôte et la mémoire du périphérique est lent. Les algorithmes doivent minimiser ces transferts en effectuant autant de calculs que possible sur le GPU.
Plusieurs modèles de conception d'algorithmes parallèles sont couramment utilisés lors du développement de noyaux GPU :
-
Mappage : 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 jeux de données.
-
Réduction : La réduction parallèle est utilisée pour calculer efficacement une seule valeur (par exemple, la somme, le maximum) à partir d'un grand jeu de données d'entrée.Voici la traduction française du fichier Markdown :
Les threads effectuent des réductions locales, qui sont ensuite combinées pour produire le résultat final.
-
Scan : Également connu sous le nom de somme préfixe, le scan est utilisé pour calculer la somme cumulée des éléments d'un tableau. Les algorithmes de scan parallèle 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 données voisines. Les calculs de stencil sont courants dans les simulations scientifiques et les applications de traitement d'images.
-
Gather/Scatter : Les threads lisent (gather) ou écrivent (scatter) à des emplacements arbitraires dans la mémoire globale. Une disposition et des schémas d'accès aux données soigneusement conçus sont nécessaires pour une efficacité optimale.
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 tenir compte avec soin de facteurs tels que l'organisation des threads, la divergence des branches, les schémas 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 et outils de programmation doivent également progresser pour permettre aux développeurs d'utiliser efficacement les nouvelles fonctionnalités et capacités du matériel. Les recherches en cours dans des domaines tels que la conception de langages de programmation, l'optimisation des compilateurs et l'autoréglage seront essentielles pour améliorer la productivité des programmeurs et la portabilité des performances à l'ère du calcul hétérogène.