Chapitre 6 : Métriques de performance GPU et analyse

L'analyse et l'optimisation des performances des applications GPU sont essentielles pour atteindre une haute efficacité et une utilisation optimale des ressources matérielles GPU. Dans ce chapitre, nous explorerons les principales métriques de performance GPU, les outils de profilage et d'optimisation, les techniques d'identification des goulots d'étranglement de performance et les stratégies d'amélioration des performances GPU.

Débit, latence et bande passante mémoire

Trois métriques fondamentales pour évaluer les performances des GPU sont le débit, la latence et la bande passante mémoire. Comprendre ces métriques et leurs implications est essentiel pour analyser et optimiser les applications GPU.

Débit

Le débit fait référence au nombre d'opérations ou de tâches qu'un GPU peut effectuer dans un temps donné. Il est généralement mesuré en opérations à virgule flottante par seconde (FLOPS) ou en instructions par seconde (IPS). Les GPU sont conçus pour atteindre un débit élevé en exploitant le parallélisme et en exécutant un grand nombre de threads simultanément.

Le débit théorique maximal d'un GPU peut être calculé à l'aide de la formule suivante :

Débit maximal (FLOPS) = Nombre de cœurs CUDA × Fréquence d'horloge × FLOPS par cœur CUDA par cycle

Par exemple, un GPU NVIDIA GeForce RTX 2080 Ti a 4352 cœurs CUDA, une fréquence d'horloge de base de 1350 MHz et chaque cœur CUDA peut effectuer 2 opérations à virgule flottante par cycle (FMA - Fused Multiply-Add). Par conséquent, son débit théorique maximal est :

Débit maximal (FLOPS) = 4352 × 1350 MHz × 2 = 11,75 TFLOPS

Cependant, atteindre le débit théorique maximal dans la pratique est difficile en raison de divers facteurs tels que les modèles d'accès à la mémoire, la divergence des branches et les contraintes de ressources.

Latence

La latence fait référence au temps nécessaire pour qu'une seule opération ou tâche soit terminée. Dans le contexte des GPU, la latence est souvent associée aux opérations d'accès à la mémoire. Les GPU ont un système mémoire hiérarchique, et l'accès aux données à différents niveaux de la hiérarchie mémoire entraîne des latences différentes.Fichier Markdown traduit en français :

Latences typiques pour différents niveaux de mémoire dans un GPU :

Registres : 0-1 cycles
Mémoire partagée : 1-2 cycles
Cache L1 : 20-30 cycles
Cache L2 : 200-300 cycles
Mémoire globale (DRAM) : 400-800 cycles

La latence peut avoir un impact significatif sur les performances du GPU, en particulier lorsqu'il y a des dépendances entre les opérations ou lorsque les threads attendent que les données soient récupérées depuis la mémoire. Des techniques comme le masquage de latence, la préfétation et la mise en cache peuvent aider à atténuer l'impact de la latence sur les performances du GPU.

Bande passante mémoire

La bande passante mémoire fait référence au débit auquel les données peuvent être transférées entre le GPU et son sous-système mémoire. Elle est généralement mesurée en octets par seconde (o/s) ou en gigaoctets par seconde (Go/s). Les GPU ont des interfaces mémoire à haute bande passante, comme GDDR6 ou HBM2, pour prendre en charge la nature gourmande en données des charges de travail graphiques et de calcul.

La bande passante mémoire maximale théorique d'un GPU peut être calculée à l'aide de la formule suivante :

Bande passante mémoire maximale (Go/s) = Fréquence d'horloge mémoire × Largeur de bus mémoire ÷ 8

Par exemple, un GPU NVIDIA GeForce RTX 2080 Ti a une fréquence d'horloge mémoire de 7000 MHz (effective) et une largeur de bus mémoire de 352 bits. Par conséquent, sa bande passante mémoire maximale théorique est :

Bande passante mémoire maximale (Go/s) = 7000 MHz × 352 bits ÷ 8 = 616 Go/s

La bande passante mémoire est un facteur critique dans les performances du GPU, car de nombreuses applications GPU sont limitées par la mémoire, ce qui signifie que leurs performances sont limitées par le débit auquel les données peuvent être transférées entre le GPU et la mémoire. Optimiser les modèles d'accès à la mémoire, minimiser les transferts de données et exploiter la hiérarchie mémoire peuvent contribuer à améliorer l'utilisation de la bande passante mémoire.

Outils de profilage et d'optimisation des performances

Les outils de profilage et d'optimisation des performances sont essentiels pour analyser le comportement des applications GPU, identifier les goulots d'étranglement des performances et guider les efforts d'optimisation. Ces outils fournissent des informations sur divers aspects des performances du GPU, tels que le temps d'exécution des noyaux, les accès mémoire, etc.Voici la traduction française du fichier Markdown :

Modèles ESS, occupation et utilisation des ressources

Parmi les outils de profilage et d'optimisation des performances populaires pour les GPU, on peut citer :

NVIDIA Visual Profiler (nvvp) : Un outil de profilage graphique qui fournit une vue d'ensemble des performances des applications GPU. Il permet aux développeurs d'analyser l'exécution des noyaux, les transferts de mémoire et les appels d'API, et fournit des recommandations pour l'optimisation.
NVIDIA Nsight : Un environnement de développement intégré (IDE) qui inclut des capacités de profilage et de débogage pour les applications GPU. Il prend en charge divers langages de programmation et frameworks, tels que CUDA, OpenCL et OpenACC.
NVIDIA Nsight Compute : Un outil de profilage autonome qui se concentre sur l'analyse des performances des noyaux GPU. Il fournit des métriques de performance détaillées, telles que le débit d'instructions, l'efficacité de la mémoire et l'occupation, et aide à identifier les goulots d'étranglement de performance au niveau du code source.
AMD Radeon GPU Profiler (RGP) : Un outil de profilage pour les GPU AMD qui capture et visualise les données de performance pour les applications DirectX, Vulkan et OpenCL. Il fournit des informations sur l'utilisation du GPU, l'utilisation de la mémoire et les blocages de pipeline.
AMD Radeon GPU Analyzer (RGA) : Un outil d'analyse statique qui analyse le code de shader GPU et fournit des prédictions de performances, l'utilisation des ressources et des suggestions d'optimisation.

Ces outils fonctionnent généralement en instrumentant le code de l'application GPU, en collectant les données de performance pendant l'exécution et en les présentant dans un format convivial pour l'analyse. Ils offrent souvent des vues chronologiques, des compteurs de performances et une corrélation avec le code source pour aider les développeurs à identifier les problèmes de performances et à optimiser leur code.

Exemple : Profilage d'une application CUDA à l'aide de NVIDIA Visual Profiler (nvvp)

Compilez l'application CUDA avec le profilage activé :
```
nvcc -o myapp myapp.cu -lineinfo
```
Exécutez l'application avec le profilage :
```
nvprof ./myapp
```
Ouvrez le Visual Profiler :
```
nvvp
```
Importez les données de profilage générées.Voici la traduction française du fichier Markdown, avec les commentaires traduits mais le code non traduit.
Analysez la vue de la chronologie, les performances des noyaux, les transferts de mémoire et les appels d'API.
Identifiez les goulots d'étranglement des performances et optimisez le code en fonction des recommandations du profileur.

Identification des goulots d'étranglement des performances

L'identification des goulots d'étranglement des performances est essentielle pour optimiser les applications GPU. Les goulots d'étranglement des performances peuvent provenir de divers facteurs, tels que des modèles d'accès à la mémoire inefficaces, une faible occupation, une divergence des branches et des contraintes de ressources. Voici quelques techniques courantes pour identifier les goulots d'étranglement des performances :

Profilage : L'utilisation d'outils de profilage pour mesurer le temps d'exécution des noyaux, le temps de transfert de mémoire et la surcharge des API peut aider à identifier les parties de l'application qui consomment le plus de temps et de ressources.
Analyse de l'occupation : L'occupation fait référence au rapport entre les warps actifs et le nombre maximum de warps pris en charge par un GPU. Une faible occupation peut indiquer une sous-utilisation des ressources GPU et peut suggérer la nécessité d'optimiser les dimensions des blocs et de la grille ou de réduire l'utilisation des registres et de la mémoire partagée.
Examen des modèles d'accès à la mémoire : Des modèles d'accès à la mémoire inefficaces, tels que des accès à la mémoire globale non regroupés ou fréquents, peuvent avoir un impact significatif sur les performances du GPU. L'analyse des modèles d'accès à la mémoire à l'aide d'outils de profilage peut aider à identifier les opportunités d'optimisation, comme l'utilisation de la mémoire partagée ou l'amélioration de la localité des données.
Étude de la divergence des branches : La divergence des branches se produit lorsque les threads d'un warp prennent des chemins d'exécution différents en raison d'instructions conditionnelles. Les branches divergentes peuvent entraîner une sérialisation et une réduction des performances. L'identification et la minimisation de la divergence des branches peuvent contribuer à améliorer les performances du GPU.
Surveillance de l'utilisation des ressources : Les GPU ont des ressources limitées, telles que les registres, la mémoire partagée et les blocs de threads. La surveillance de l'utilisation des ressources à l'aide d'outils de profilage peut aider à identifier les goulots d'étranglement liés aux ressources et à guider les efforts d'optimisation, comme la réduction de l'utilisation des registres.Voici la traduction française du fichier Markdown, avec les commentaires traduits mais le code non traduit :

Exemple : Identifier un goulot d'étranglement d'accès à la mémoire à l'aide de NVIDIA Nsight Compute

Profilez l'application CUDA à l'aide de Nsight Compute :
```
ncu -o profile.ncu-rep ./myapp
```
Ouvrez le rapport de profilage généré dans Nsight Compute.
Analysez la section "Analyse de la charge de travail mémoire" pour identifier les modèles d'accès mémoire inefficaces, tels que les accès non regroupés ou l'utilisation élevée de la mémoire globale.
Optimisez les modèles d'accès mémoire en fonction des informations fournies par Nsight Compute, comme l'utilisation de la mémoire partagée ou l'amélioration de la localité des données.

Stratégies pour améliorer les performances du GPU

Une fois que les goulots d'étranglement de performances ont été identifiés, diverses stratégies peuvent être employées pour améliorer les performances du GPU. Voici quelques stratégies d'optimisation courantes :

Maximiser le parallélisme : Assurez-vous que l'application est décomposée en un nombre suffisant de tâches parallèles pour utiliser pleinement les ressources du GPU. Cela peut impliquer d'ajuster les dimensions des blocs et de la grille, d'utiliser des flux pour une exécution concurrente ou d'exploiter le parallélisme au niveau des tâches.
Optimiser les modèles d'accès mémoire : Améliorez l'efficacité des accès mémoire en minimisant les accès à la mémoire globale, en utilisant la mémoire partagée pour les données fréquemment accédées et en assurant des accès mémoire regroupés. Des techniques telles que le découpage de la mémoire, les transformations de la disposition des données et la mise en cache peuvent aider à optimiser les performances de la mémoire.
Réduire la divergence des branches : Minimisez la divergence des branches en restructurant le code pour éviter les branches divergentes au sein d'un warp. Des techniques comme la prédiction de branches, les branches dépendantes des données et la programmation au niveau des warps peuvent aider à réduire l'impact de la divergence des branches.
Exploiter la hiérarchie de la mémoire : Tirez parti efficacement de la hiérarchie de la mémoire du GPU en maximisant l'utilisation des registres et de la mémoire partagée pour les données fréquemment accédées. Utilisez la mémoire de texture et la mémoire constante pour les données en lecture seule qui présentent une localité spatiale ou sont accédées de manière uniforme entre les threads.
Chevaucher les calculs et les accès mémoire : Exploitez le parallélisme entre les calculs et les accès mémoire pour masquer les latences mémoire. Utilisez des techniques telles que le lancement anticipé de noyaux, l'utilisation de streams et l'asynchronisme des transferts de données pour optimiser l'utilisation du GPU.Voici la traduction française du fichier Markdown, avec les commentaires traduits mais le code non traduit :

Optimisation des performances GPU

Utilisation de la mémoire partagée : Tirez parti de la mémoire partagée (shared memory) pour réduire les accès à la mémoire globale (global memory), qui sont plus lents. Stockez les données fréquemment utilisées dans la mémoire partagée pour améliorer les performances.
Coalescing des accès mémoire : Assurez-vous que les threads d'un même warp accèdent à des emplacements mémoire adjacents. Cela permet d'optimiser les accès à la mémoire globale et d'améliorer les performances.
Éviter les divergences de branchement : Minimisez les divergences de branchement (branch divergence) au sein d'un même warp, car elles peuvent entraîner une exécution séquentielle des branches, réduisant ainsi les performances.
Overlapping des transferts mémoire : Masquez la latence des transferts mémoire en chevauchant les calculs avec les transferts de mémoire à l'aide de flux CUDA (CUDA streams) ou de files de commandes OpenCL (OpenCL command queues). Cela permet au GPU d'effectuer des calculs pendant que les données sont transférées entre la mémoire hôte et la mémoire du périphérique.
Réglage des paramètres de lancement des noyaux : Expérimentez avec différentes tailles de blocs et de grilles pour trouver la configuration optimale pour chaque noyau. Les paramètres de lancement optimaux dépendent de facteurs tels que le nombre de registres utilisés par thread, l'utilisation de la mémoire partagée et les caractéristiques de l'architecture GPU.
Minimisation des transferts de données hôte-périphérique : Réduisez la quantité de données transférées entre l'hôte (CPU) et le périphérique (GPU) en effectuant le plus de calculs possible sur le GPU. Regroupez les petits transferts en transferts plus importants pour amortir le surcoût de chaque transfert.
Utilisation d'opérations asynchrones : Tirez parti des opérations asynchrones, telles que les copies mémoire asynchrones et les lancements de noyaux asynchrones, pour chevaucher les calculs et les communications. Cela permet au CPU d'effectuer d'autres tâches pendant que le GPU exécute, améliorant ainsi les performances globales de l'application.

Exemple : Optimisation des schémas d'accès mémoire à l'aide de la mémoire partagée en CUDA

Code d'origine avec des accès inefficaces à la mémoire globale :

__global__ void myKernel(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < n; i++) {
            result += data[tid] * data[i];
        }
        data[tid] = result;
    }
}

Code optimisé utilisant la mémoire partagée :

__global__ void myKernel(float* data, int n) {
    __shared__ float sharedData[256];
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int localIdx = threadIdx.x;
 
    if (tid < n) {
        sharedData[localIdx] = data[tid];
    }
    __syncthreads();
 
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < blockDim.x; i++) {
            result += sharedData[localIdx] * sharedData[i];
        }
        data[tid] = result;
    }
}
```Voici la traduction française du fichier Markdown, avec les commentaires traduits mais le code non traduit :
 
a[tid] = result;
    }
}

Dans le code optimisé, les données d'entrée sont d'abord chargées dans la mémoire partagée, qui a une latence beaucoup plus faible par rapport à la mémoire globale. Le calcul est ensuite effectué en utilisant la mémoire partagée, réduisant ainsi le nombre d'accès à la mémoire globale et améliorant les performances.

Conclusion

L'analyse et l'optimisation des performances des GPU sont essentielles pour développer des applications GPU efficaces et performantes. En comprenant les principales métriques de performance telles que le débit, la latence et la bande passante mémoire, les développeurs peuvent prendre des décisions éclairées sur l'optimisation de leur code.

Les outils de profilage et d'optimisation des performances jouent un rôle crucial dans l'identification des goulots d'étranglement des performances et guident les efforts d'optimisation. Ces outils fournissent des informations précieuses sur l'exécution des noyaux, les modèles d'accès à la mémoire, l'occupation et l'utilisation des ressources, permettant aux développeurs de concentrer leurs efforts d'optimisation sur les domaines les plus critiques.

Voici quelques stratégies courantes pour optimiser les performances des GPU, poursuivies au format Markdown :

Réduction de 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 réduction de l'efficacité SIMD. Les algorithmes doivent être structurés pour minimiser la divergence des branches dans la mesure du possible. Des techniques telles que la prédiction de branches, les branches dépendantes des données et la programmation au niveau des warps peuvent aider à réduire l'impact de la divergence des branches.
Exploitation de la hiérarchie de la mémoire : Exploitez efficacement la hiérarchie de la mémoire GPU en maximisant l'utilisation des registres et de la mémoire partagée pour les données fréquemment accédées. Utilisez la mémoire texture et la mémoire constante pour les données en lecture seule qui présentent une localité spatiale ou sont accédées de manière uniforme entre les threads.
Chevauchement du calcul et des transferts de mémoire : Masquez la latence des transferts de mémoire en chevauchant le calcul avec les transferts de mémoire à l'aide de flux CUDA ou de files de commandes OpenCL. Cela permetVoici la traduction française du fichier Markdown, avec les commentaires traduits mais le code non traduit :
Ajuster les paramètres de lancement du noyau : Expérimentez avec différentes tailles de blocs et de grilles pour trouver la configuration optimale pour chaque noyau. Les paramètres de lancement optimaux dépendent de facteurs tels que le nombre de registres utilisés par thread, l'utilisation de la mémoire partagée et les caractéristiques de l'architecture GPU.
Minimiser les transferts de données hôte-périphérique : Réduisez la quantité de données transférées entre l'hôte (CPU) et le périphérique (GPU) en effectuant le plus de calculs possible sur le GPU. Regroupez les petits transferts en transferts plus importants pour amortir le surcoût de chaque transfert.
Utiliser des opérations asynchrones : Tirez parti des opérations asynchrones, telles que les copies de mémoire asynchrones et les lancements de noyaux asynchrones, pour chevaucher les calculs et les communications. Cela permet au CPU d'effectuer d'autres tâches pendant que le GPU exécute, améliorant ainsi les performances globales de l'application.

Exemple : Optimisation des schémas d'accès à la mémoire à l'aide de la mémoire partagée dans CUDA

Code d'origine avec des accès à la mémoire globale inefficaces :

__global__ void myKernel(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < n; i++) {
            result += data[tid] * data[i];
        }
        data[tid] = result;
    }
}

Code optimisé à l'aide de la mémoire partagée :

__global__ void myKernel(float* data, int n) {
    __shared__ float sharedData[256];
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int localIdx = threadIdx.x;
 
    if (tid < n) {
        sharedData[localIdx] = data[tid];
    }
    __syncthreads();
 
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < blockDim.x; i++) {
            result += sharedData[localIdx] * sharedData[i];
        }
        data[tid] = result;
    }
}

Dans le code optimisé, les données d'entrée sont d'abord chargées dans la mémoire partagée, qui a une latence beaucoup plus faible par rapport àVoici la traduction française du fichier Markdown, avec les commentaires traduits mais le code laissé tel quel :

Mémoire globale. Le calcul est ensuite effectué à l'aide de la mémoire partagée, réduisant le nombre d'accès à la mémoire globale et améliorant les performances.

// Allouer de la mémoire globale pour les données d'entrée et de sortie
float *d_input, *d_output;
cudaMalloc(&d_input, size * sizeof(float));
cudaMalloc(&d_output, size * sizeof(float));
 
// Copier les données d'entrée depuis l'hôte vers le périphérique
cudaMemcpy(d_input, h_input, size * sizeof(float), cudaMemcpyHostToDevice);
 
// Lancer le kernel
dim3 block(BLOCK_SIZE, BLOCK_SIZE);
dim3 grid(size / block.x, size / block.y);
kernel<<<grid, block>>>(d_input, d_output, size);
 
// Copier les résultats depuis le périphérique vers l'hôte
cudaMemcpy(h_output, d_output, size * sizeof(float), cudaMemcpyDeviceToHost);
 
// Libérer la mémoire allouée
cudaFree(d_input);
cudaFree(d_output);

Chapter 5 Gpu Memory System Design Chapter 7 Streaming Multiprocessor Design