Comment concevoir des puces GPU
Chapter 7 Streaming Multiprocessor Design

Chapitre 7 : Conception du processeur de flux multiprocesseur dans la conception GPU

Le processeur de flux multiprocesseur (SM) est le bloc de construction fondamental des architectures GPU NVIDIA. Chaque SM contient un ensemble de cœurs CUDA qui exécutent les instructions dans un mode SIMT (Single Instruction, Multiple Thread). Le SM est responsable de la gestion et de l'ordonnancement des warps, de la gestion de la divergence des branches et de l'accès rapide à la mémoire partagée et aux caches. Dans ce chapitre, nous explorerons la micro-architecture du SM, y compris ses pipelines, ses mécanismes d'ordonnancement des warps, la conception de son registre et l'organisation de sa mémoire partagée et de son cache L1.

Micro-architecture et pipelines du SM

Le SM est un processeur très parallèle et pipeliné conçu pour exécuter efficacement des centaines de threads simultanément. La figure 7.1 montre un schéma bloc simplifié d'un SM dans l'architecture NVIDIA Volta.

                                 Cache d'instructions
                                         |
                                         v
                                    Ordonnanceur de warps
                                         |
                                         v
                               Unité de répartition (4 warps)
                                 |   |   |   |
                                 v   v   v   v
                               Cœur CUDA (FP64/FP32/ENT)
                               Cœur CUDA (FP64/FP32/ENT)
                               Cœur CUDA (FP64/FP32/ENT)
                               ...
                               Tenseur Core
                               Tenseur Core
                               ...
                               Unité de chargement/stockage
                               Unité de chargement/stockage
                               ...
                               Unité de fonctions spéciales
                                         ^
                                         |
                                Registre (64 Ko)
                                         ^
```Voici la traduction française du fichier Markdown avec les commentaires traduits, mais le code non traduit :

                                  Mémoire partagée / Cache L1 (96 Ko)

Figure 7.1 : Diagramme de bloc simplifié d'un SM dans l'architecture NVIDIA Volta.

Les principaux composants du SM comprennent :

  1. Cache d'instructions : Stocke les instructions fréquemment accédées afin de réduire la latence et d'améliorer le débit.

  2. Ordonnanceur de warps : Sélectionne les warps prêts à s'exécuter et les dispatche vers les unités d'exécution disponibles.

  3. Unité de dispatch : Récupère et décode les instructions pour jusqu'à 4 warps par cycle et les dispatche vers les unités d'exécution appropriées.

  4. Cœurs CUDA : Unités d'exécution programmables qui prennent en charge une large gamme d'opérations entières et à virgule flottante. Chaque SM de Volta contient 64 cœurs CUDA.

  5. Tensor Cores : Unités d'exécution spécialisées conçues pour accélérer les charges de travail d'apprentissage en profondeur et d'IA. Chaque SM de Volta contient 8 Tensor Cores.

  6. Unités de chargement/stockage : Gèrent les opérations mémoire, y compris les chargements et les stockages en mémoire globale, en mémoire partagée et dans les caches.

  7. Unités de fonctions spéciales : Exécutent les opérations transcendantales et autres opérations mathématiques complexes.

  8. Registre de fichiers : Fournit un accès rapide aux registres privés des threads. Chaque SM de Volta dispose d'un registre de fichiers de 64 Ko.

  9. Mémoire partagée / Cache L1 : Un espace mémoire configurable qui peut être utilisé comme un cache géré par le logiciel (mémoire partagée) ou comme un cache de données L1 géré par le matériel.

Le pipeline SM est conçu pour maximiser le débit en permettant l'exécution concurrente de plusieurs warps et en masquant la latence mémoire. La figure 7.2 illustre une vue simplifiée du pipeline SM.

    Récupération d'instructions
            |
            v
    Décodage d'instructions
            |
            v
    Collecte d'opérandes
            |
            v
    Exécution (Cœurs CUDA, Tensor Cores, Unités de chargement/stockage, Unités de fonctions spéciales)
            |
            v
    Écriture des résultats

Figure 7.2 : Pipeline SM simplifié.

Les étapes du pipeline sont les suivantes :

  1. Récupération d'instructions : L'ordonnanceur de warps sélectionne un warp prêt à s'exécuter et récupère ses instructions.Voici la traduction française :

  2. Instruction Fetch: Le processeur stream (SM) récupère l'instruction suivante pour ce warp à partir du cache d'instructions.

  3. Décodage d'instruction: L'instruction récupérée est décodée pour déterminer le type d'opération, les opérandes et les registres de destination.

  4. Collecte des opérandes: Les opérandes requises pour l'instruction sont collectées à partir du registre ou de la mémoire partagée.

  5. Exécution: L'instruction est exécutée sur l'unité d'exécution appropriée (cœur CUDA, cœur tensoriel, unité de chargement/stockage ou unité de fonctions spéciales).

  6. Écriture de retour: Le résultat de l'exécution est écrit de retour dans le registre ou la mémoire partagée.

Pour atteindre de hautes performances, le SM utilise plusieurs techniques pour maximiser l'utilisation des ressources et masquer la latence :

  • Émission double: Le SM peut émettre deux instructions indépendantes par warp en un seul cycle, permettant d'augmenter le parallélisme au niveau de l'instruction.
  • Unités d'exécution pipelinées: Les unités d'exécution sont pipelinées, permettant au SM de démarrer une nouvelle opération sur une unité avant que l'opération précédente ne soit terminée.
  • Masquage de la latence: Le SM peut basculer entre les warps sur une base cycle par cycle, lui permettant de masquer la latence des accès mémoire et des opérations à longue latence en exécutant des instructions d'autres warps.

L'exemple 7.1 montre un noyau CUDA simple qui effectue l'addition élément par élément de deux vecteurs.

__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

Exemple 7.1 : Noyau CUDA pour l'addition de vecteurs.

Dans cet exemple, chaque thread du noyau calcule la somme des éléments correspondants des vecteurs d'entrée a et b et stocke le résultat dans le vecteur de sortie c. Le SM exécute ce noyau en affectant chaque thread à un cœur CUDA et en ordonnançant des warps de threads pour les exécuter sur les cœurs disponibles. Les unités de chargement/stockage sont utilisées pour récupérer les données d'entrée depuis la mémoire globale et écrire les résultats.

Ordonnancement des warps et gestion des divergences

EfVoici la traduction française du fichier markdown, avec les commentaires traduits mais le code laissé inchangé :

Un ordonnancement efficace des warps est crucial pour maximiser les performances du SM. L'ordonnanceur de warps est responsable de la sélection des warps prêts à s'exécuter et de leur envoi aux unités d'exécution disponibles. L'objectif principal de l'ordonnanceur de warps est de maintenir les unités d'exécution occupées en s'assurant qu'il y a toujours des warps disponibles pour s'exécuter.

Le SM utilise un mécanisme d'ordonnancement de warps à deux niveaux :

  1. Ordonnancement des warps : L'ordonnanceur de warps sélectionne les warps prêts à s'exécuter en fonction d'une politique d'ordonnancement, comme le tourniquet (round-robin) ou l'ordre d'ancienneté (oldest-first). Les warps sélectionnés sont alors envoyés aux unités d'exécution disponibles.

  2. Ordonnancement des instructions : Au sein de chaque warp, le SM planifie les instructions en fonction de leurs dépendances et de la disponibilité des unités d'exécution. Le SM peut émettre plusieurs instructions indépendantes du même warp en un seul cycle pour maximiser le parallélisme au niveau des instructions.

La figure 7.3 illustre le mécanisme d'ordonnancement de warps à deux niveaux.

    Pool de warps
    Warp 1 (Prêt)
    Warp 2 (En attente)
    Warp 3 (Prêt)
    ...
    Warp N (Prêt)
        |
        v
    Ordonnanceur de warps
        |
        v
    Unité de répartition
        |
        v
    Unités d'exécution

Figure 7.3: Mécanisme d'ordonnancement de warps à deux niveaux.

L'un des défis clés de l'ordonnancement des warps est de gérer la divergence des branches. Dans le modèle d'exécution SIMT, tous les threads d'un warp exécutent la même instruction en synchronisme. Cependant, lorsqu'un warp rencontre une instruction de branchement (par exemple, une déclaration if-else), certains threads peuvent emprunter le chemin du if tandis que d'autres empruntent le chemin du else. Cette situation est appelée divergence de branche.

Pour gérer la divergence des branches, le SM utilise une technique appelée prédication. Lorsqu'un warp rencontre une branche divergente, le SM exécute les deux chemins de la branche séquentiellement, en masquant les threads qui ne prennent pas chaque chemin. Les résultats sont ensuite combinés à l'aide de registres de prédicat pour s'assurer que chaque thread reçoit le résultat correct.

L'exemple 7.2 montre un noyau CUDA avec une branche divergente.Voici la traduction française du fichier Markdown avec la traduction des commentaires pour le code, mais sans traduire le code lui-même :

__global__ void divergentKernel(int *data, int *result) {
    // Identifiant unique du thread
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (data[tid] > 0) {
        // Si la valeur de données est positive, multiplier par 2
        result[tid] = data[tid] * 2;
    } else {
        // Sinon, multiplier par 3
        result[tid] = data[tid] * 3;
    }
}

Exemple 7.2 : Noyau CUDA avec une branche divergente.

Dans cet exemple, la condition de branche data[tid] > 0 peut faire en sorte que certains threads d'un warp suivent le chemin du if, tandis que d'autres suivent le chemin du else. Le SM (Streaming Multiprocessor) gère cette divergence en exécutant les deux chemins de manière séquentielle et en masquant les threads inactifs dans chaque chemin.

La Figure 7.4 illustre le processus de prédication pour un warp avec des threads divergents.

    Warp (32 threads)
    Thread 1 : data[1] = 5, result[1] = 10
    Thread 2 : data[2] = -3, result[2] = -9
    ...
    Thread 32 : data[32] = 7, result[32] = 14

    Branche divergente :
    if (data[tid] > 0) {
        result[tid] = data[tid] * 2 ;
    } else {
        result[tid] = data[tid] * 3 ;
    }

    Prédication :
    Étape 1 : Exécuter le chemin du `if` avec un masque
        Thread 1 : result[1] = 10
        Thread 2 : (masqué)
        ...
        Thread 32 : result[32] = 14

    Étape 2 : Exécuter le chemin du `else` avec un masque
        Thread 1 : (masqué)
        Thread 2 : result[2] = -9
        ...
        Thread 32 : (masqué)

    Résultat final :
    Thread 1 : result[1] = 10
    Thread 2 : result[2] = -9
    ...
    Thread 32 : result[32] = 14

Figure 7.4 : Processus de prédication pour un warp avec des threads divergents.

En utilisant la prédication, le SM peut gérer la divergence des branches sans avoir besoin d'instructions de branchement explicites ou de divergence du flux de contrôle. Cependant, les branches divergentes peuvent toujours avoir un impact sur les performances, car le SM doit exécuter les deux chemins de manière séquentielle, réduisant ainsi le parallélisme effectif.

Registre et collecteurs d'opérandes

Le registre est un composant essentiel du SM, fournissant un accès rapide aux registres privés des threads. Chaque SM dispose d'un grand registre pour prendre en charge de nombreux threads actifs et permettre un changement de contexte efficace entre les warps.Voici la traduction française du fichier Markdown avec les commentaires traduits, mais sans traduction du code:

Dans l'architecture NVIDIA Volta, chaque SM a un fichier de registres de 64 Ko, organisé en 32 banques de 2 Ko chacune. Le fichier de registres est conçu pour fournir un débit élevé et un accès à faible latence afin de prendre en charge le grand nombre de fils d'exécution concurrents.

Pour minimiser les conflits de banque et améliorer les performances, le SM utilise une technique appelée collecte d'opérandes. Les collecteurs d'opérandes sont des unités spécialisées qui rassemblent les opérandes des banques de fichiers de registres et les transmettent aux unités d'exécution. En utilisant les collecteurs d'opérandes, le SM peut réduire l'impact des conflits de banque et améliorer l'utilisation des unités d'exécution.

La figure 7.5 montre un schéma simplifié du fichier de registres et des collecteurs d'opérandes dans un SM.

    Fichier de registres (64 Ko)
    Banque 1 (2 Ko)
    Banque 2 (2 Ko)
    ...
    Banque 32 (2 Ko)
        |
        v
    Collecteurs d'opérandes
        |
        v
    Unités d'exécution

Figure 7.5 : Fichier de registres et collecteurs d'opérandes dans un SM.

Les collecteurs d'opérandes fonctionnent en rassemblant les opérandes de plusieurs instructions et plusieurs warps, permettant au SM d'émettre des instructions de différents warps vers les unités d'exécution en un seul cycle. Cela permet de masquer la latence des accès au fichier de registres et d'améliorer le débit global du SM.

L'exemple 7.3 montre un noyau CUDA qui effectue un produit scalaire de deux vecteurs.

__global__ void dotProduct(float *a, float *b, float *result, int n) {
    // Chaque thread calcule une somme partielle du produit scalaire en utilisant son index assigné
    __shared__ float partialSum[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    partialSum[tid] = 0;
 
    while (i < n) {
        partialSum[tid] += a[i] * b[i];
        i += blockDim.x * gridDim.x;
    }
 
    __syncthreads();
 
    // Réduction de la somme partielle dans le tableau partialSum
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            partialSum[tid] += partialSum[tid + s];
        }
        __syncthreads();
    }
 
    if (tid == 0) {
        result[blockIdx.x] = partialSum[0];
    }
}

Dans cet exemple, chaque thread calcule une somme partielle du produit scalaire en utilisant son index assigné.Voici la traduction française du fichier markdown :

Éléments des vecteurs d'entrée. Les sommes partielles sont stockées dans le tableau de mémoire partagée partialSum. Après que tous les threads aient calculé leurs sommes partielles, une réduction parallèle est effectuée pour additionner les sommes partielles et obtenir le résultat final du produit scalaire.

Le collecteur d'opérandes joue un rôle crucial dans cet exemple en rassemblant efficacement les opérandes pour les accès à la mémoire partagée et les opérations arithmétiques. Il aide à éviter les conflits de banque et améliore l'utilisation des unités d'exécution.

Conclusion

Le multiprocesseur de flux (streaming multiprocessor) est l'unité de calcul de base dans les architectures GPU modernes. Sa conception se concentre sur la maximisation du débit et la dissimulation de la latence de la mémoire grâce à une combinaison de multithreading à grain fin, d'exécution SIMT et de collecte efficace des opérandes.

Les composants clés du SM incluent l'ordonnanceur de warps, qui sélectionne les warps pour l'exécution ; la pile SIMT, qui gère la divergence et la convergence des branches ; le fichier de registres et les collecteurs d'opérandes, qui fournissent un accès rapide aux registres privés des threads ; et la mémoire partagée et le cache L1, qui permettent un partage et une réutilisation des données à faible latence.

À mesure que les architectures GPU continuent d'évoluer, la recherche dans des domaines tels que la gestion de la divergence des branches, l'ordonnancement des warps et la conception du fichier de registres sera cruciale pour améliorer les performances et l'efficacité des futurs GPU. Des techniques novatrices comme la formation dynamique de warps, la compaction de blocs de threads et les caches de réutilisation des opérandes ont le potentiel d'améliorer de manière significative les capacités du SM et d'atteindre de nouveaux niveaux de performance dans les charges de travail de calcul parallèle.

Chapter 6 Gpu Performance Metrics and AnalysisChapter 8 Interconnect and on Chip Networks