\chapter{Architectures parall\`eles de GPUs} \label{chap:archi_parallel} %Introduction The high performance computing (HPC) industry’s need for computation is increasing, as large and complex computational problems become commonplace across many industry segments. Traditional CPU technology, however, is no longer capable of scaling in performance sufficiently to address this demand. The parallel processing capability of the GPU allows it to divide complex computing tasks into thousands of smaller tasks that can be run concurrently. This ability is enabling computational scientists and researchers to address some of the world’s most challenging computational problems up to several orders of magnitude faster. At the end of 2006, NVIDIA GPUs have invaded the field of general purpose computing with the invention of Compute Unified Device Architecture (CUDA). The performance and memory bandwidth of the NVIDIA GPUs make them an attractive choice for solving many scientific computational problems. In November 2006, NVIDIA introduced CUDATM, a general purpose parallel computing architecture – with a new parallel programming model and instruction set architecture – that leverages the parallel compute engine in NVIDIA GPUs to solve many complex computational problems in a more efficient way than on a CPU. L'architecture GPU utilisée dans ce document est celle basée sur la plateforme CUDA (Compute Unified Device Architecture), développée par NVIDIA pour une programmation générique des GPUs afin de bénéficier de leur capacité de traitement massivement parallèle. Pourquoi parler des accelerateurs graphiques ? Ce cours propose de decouvrir et/ou de comprendre les raisons du succes des accelerateurs de calcul dans le monde du calcul haute performance (HPC). Depuis plusieurs annees maintenant les constructeurs esperaient proposer des architectures materielles simples et puissantes permettant de depasser les limites atteintes par les CPU. La technologie de GPGPU ( general purpose processing on graphics processing unit ) s'est imposee au milieu d'autres technologies existantes pour des raisons d'architecture materielle mais aussi logicielle. L'objectif double de cours est de presenter ces elements d'architecture ainsi que l'adequation avec les problematiques actuelles des thematiques, gourmandes en puissance de calcul, du monde du HPC. l'un des plus grands fournisseurs de processeurs graphiques: NVIDIA définition HPC: high performance computing %%-------------------------------------------------------------------------------------------------------%% %%-------------------------------------------------------------------------------------------------------%% \section{Calcul parallèle} \label{sec:parallel} Pour définir le principe d'un calcul parallèle, nous préférons définir, tout d'abord, celui de son opposé: le calcul séquentiel, pour bien cerner la différence entre ces deux manières de calcul. Un calcul séquentiel consiste en l'exécution d'un programme, instruction par instruction, par un seul processeur (unité de calcul) et de façon à ce qu'une seule instruction soit exécutée à la fois. En revanche, un calcul parallèle est défini comme l'exécution d'un même programme, simultanément, par plusieurs processeurs. Nous avons, en général, deux façons de réaliser un calcul parallèle. La première consiste en le découpage du programme en plusieurs tâches de calcul puis, d'exécuter toutes ces tâches en parallèle par différents processeurs. La seconde nécessite d'abord le partitionnement des données du problème à traiter, de manière à ce que chaque partie de données soit attribuée à un processeur différent. Ensuite, tous les processeurs exécutent en parallèle les mêmes instructions du programme mais, opérant sur des données différentes. Cette dernière méthode, appelée la \textit{parallélisation des données}, est celle retenue dans ce document. En outre, les calculs parallèles nécessitent aussi une gestion des dépendances de données entre les différents processeurs. Les calculs locaux de deux processeurs sont dits dépendants lorsque l'exécution de l'un affecte le résultat de l'autre. Une dépendance de données implique une utilisation de la valeur d'une même variable par les calculs locaux de deux ou plusieurs processeurs. Les dépendances de données peuvent être gérées par la synchronisation des lectures/écritures dans une même mémoire (systèmes à mémoires partagées) ou par la communication de données entre processeurs via des messages (systèmes à mémoires distribuées). Le calcul parallèle a pour objectif d'exploiter la grande quantité de ressources (puissance de calcul, espace mémoire, ...) que permettent d'offrir les calculateurs parallèles, pour réduire le temps d'exécution des applications nécessitant un long traitement et/ou pour pouvoir exécuter celles portant sur des volumes de données très importants. Ainsi, il nous permet d'aborder de nouveaux problèmes, de plus en plus, complexes et de tailles toujours croissantes. \subsection{Classification des architectures parallèles} Un calculateur parallèle peut être, tout simplement, un processeur multic\oe ures possédant au moins deux unités de calcul physiques gravées sur la même puce, un supercalculateur qui permet de rassembler les composantes de plusieurs ordinateurs (processeurs et mémoires) dans une seule machine ou, une plateforme distribuée composée de plusieurs ordinateurs indépendants, homogènes ou hétérogènes, reliés entre eux par un réseau de communication. Il existe dans la littérature plusieurs classifications pour les architectures des calculateurs parallèles, basées sur différents critères de classification~\cite{ref21,ref22,ref23,ref24}. Dans cette section, nous présentons la classification la plus largement utilisée dans le domaine du calcul parallèle, nommée la \textit{taxonomie de Flynn}~\cite{ref21}. Elle est basée sur deux critères: le nombre d'instructions et le nombre de données, qui peuvent être traitées, simultanément, par les différents processeurs du calculateur parallèle. Les quatre catégories possibles de la taxonomie de Flynn sont les suivantes. \subsubsection{Instruction unique, donnée unique (SISD)} La classe SISD (Single Instruction, Single Data) représente l'ensemble des calculateurs séquentiels à une seule unité de calcul (ou monoprocesseur). Ce sont les calculateurs qui ne sont capables de traiter qu'une seule instruction sur une seule donnée, par cycle d'horloge. Bien évidemment, cette catégorie n'est pas une architecture parallèle. \subsubsection{Instructions multiples, donnée unique (MISD)} La classe MISD (Multiple Instruction, Single Data) correspond aux calculateurs parallèles pouvant exécuter plusieurs instructions, simultanément, sur la même donnée. Peu de calculateurs MISD ont existé en pratique, vu le nombre réduit des applications qui peuvent être mises en \oe uvre sur ce type d'architecture. Un exemple de calculateur parallèle expérimental MISD a été développé à l'université de Carnegie Mellon~\cite{ref25}. \subsubsection{Instruction unique, données multiples (SIMD)} La classe SIMD (Single Instruction, Multiple Data) correspond aux processeurs vectoriels et, plus généralement, aux calculateurs composés d'un grand nombre d'unités de calcul. Chaque processeur d'un calculateur SIMD exécute la même instruction à chaque cycle d'horloge, mais opérant sur des données différentes. Cette architecture parallèle est bien adaptée aux traitement des problèmes à structure régulière, où la même instruction est appliquée à un ensemble de données (exécutions des opérations sur des vecteurs ou des tableaux). \subsubsection{Instructions multiples, données multiples (MIMD)} La classe MIMD (Multiple Instruction, Multiple Data) représente la classe la plus générale dans cette classification. Les calculateurs parallèles MIMD possèdent plusieurs processeurs interconnectés entre eux, tels que chaque processeur est capable de suivre son propre chemin d'exécution. En effet, à chaque cycle d'horloge, les processeurs peuvent exécuter des instructions différentes sur des données différentes. \subsection{Mémoires des architectures parallèles} Nous pouvons distinguer, en général, deux modèles de gestion de la mémoire des calculateurs parallèles. \subsubsection{Mémoire partagée} Dans ce type d'architecture, \subsubsection{Mémoire distribuée} %%-------------------------------------------------------------------------------------------------------%% %%-------------------------------------------------------------------------------------------------------%% \section{Unité de traitement graphique GPU} \label{sec:GPU} L'architecture et l'environnement de programmation des GPUs utilisés dans ce document sont ceux basés sur la plateforme CUDA (Compute Unified Device Architecture) développée par NVIDIA~\cite{ref19}. \subsection{Architecture matérielle GPU} \label{sec:archiGPU} Les processeurs graphiques GPUs sont initialement conçus pour le traitement des applications graphiques et de la visualisation 3D. Nous pouvons citer, par exemple, les produits \textit{GeForce} et \textit{Quadro}, deux gammes de GPUs proposées par NVIDIA, qui sont destinés, respectivement, au graphisme grand public et à la visualisation professionnelle. Depuis quelques années, les GPUs sont devenus des outils très attrayants pour le calcul haute performance (HPC). La gamme de produits \textit{Tesla} a été conçue par NVIDIA pour offrir des capacités de calcul parallèle élevées et assister les processeurs dans les calculs intensifs des applications scientifiques et/ou industrielles. La figure~\ref{fig:archi} montre les différentes architectures matérielles GPU développées par NVIDIA. \begin{figure}[!h] \centering \includegraphics[width=85mm,keepaspectratio]{Figures/archi} \caption{Historique des architectures matérielles GPU.} \label{fig:archi} \end{figure} Un GPU est un processeur graphique relié à un processeur traditionnel (CPU) via un PCI-Express (voir figure~\ref{fig:gpu-cpu}). Il est souvent considéré comme un accélérateur des tâches parallèles et des opérations arithmétiques intensives d'une application exécutée sur un CPU. Il puise sa puissance de calcul de son architecture matérielle et logicielle massivement parallèle. En effet, à la différence d'une architecture CPU, un GPU est composé de centaines (voire de milliers) de processeurs (SP), appelés communément c\oe urs, organisés en plusieurs blocs de processeurs appelés multiprocesseurs (SM ou SMX). La figure~\ref{fig:compar} montre une comparaison entre l'architecture matérielle d'un CPU et celle d'un GPU Fermi. Les processeurs d'un GPU sont regroupés par 8 (Tesla), 32 (Fermi) ou 192 (Kepler) dans un multiprocesseur, selon le type de son architecture matérielle. De la même manière, les multiprocesseurs sont eux-mêmes regroupés par 2 (G80) ou 3 (GT200) dans un TPC (Texture Processing Cluster) pour l'architecture Tesla et par 4 (Fermi) ou 2 (Kepler) dans un GPC (Graphics Processing Cluster) pour les nouvelles architectures. \begin{figure}[!h] \centering \begin{tabular}{ccc} \includegraphics[width=45mm,keepaspectratio]{Figures/fig} & \includegraphics[width=50mm,keepaspectratio]{Figures/fig1} & \includegraphics[width=60mm,keepaspectratio]{Figures/schema1} \\ (a) Carte graphique GPU & (b) Un GPU relié à un CPU & (c) Un schéma de GPU relié à un CPU\\ \end{tabular} \caption{Un exemple de CPU équipé d'un GPU} \label{fig:gpu-cpu} \end{figure} \begin{figure}[!h] \centering \begin{tabular}{cc} \includegraphics[width=50mm,keepaspectratio]{Figures/CPU_scheme} & \includegraphics[width=50mm,keepaspectratio]{Figures/GPU_scheme} \\ (a) Un CPU à 8 c\oe urs & (b) Un GPU Fermi à 512 c\oe urs \end{tabular} \caption{Comparaison du nombre de c\oe urs dans un CPU et dans un GPU.} \label{fig:compar} \end{figure} En plus de la hiérarchie de processeurs, un GPU est doté d'une hiérarchie de mémoires de différentes tailles et de différentes bandes passantes mémoires. Nous distinguons, au total, six mémoires différentes (voir figure~\ref{fig:memoires}): \begin{figure}[!h] \centering \begin{tabular}{cc} \includegraphics[width=82mm,keepaspectratio]{Figures/memoires} & \includegraphics[width=82mm,keepaspectratio]{Figures/memoiresFermi} \\ (a) Architecture Tesla & (b) Architecture Fermi ou Kepler \\ \end{tabular} \caption{Hiérarchie de mémoires GPU.} \label{fig:memoires} \end{figure} \begin{itemize*} \item \textit{Registres}: chaque multiprocesseur a 8K à 65K de registres à 32-bit, répartis entre tous ses processeurs. Ce sont des mémoires rapides, accessibles en lecture/écriture et avec une faible latence (environ 1 cycle); \\ \item \textit{Mémoire partagée}: de 16 Ko à 48 Ko de mémoire par multiprocesseur. C'est une petite mémoire extrêmement rapide. Elle est dotée d'une large bande passante mémoire (plus d'un To/s) et d'une faible latence (environ 1 à 2 cycles). Elle est accessible en lecture/écriture par tous les processeurs du même multiprocesseur; \\ \item \textit{Mémoire globale}: chaque GPU est équipé de sa propre RAM (GDDR3 ou GDDR5) de 1 Go à 6 Go. C'est une mémoire accessible en lecture/écriture et partagée entre tous les multiprocesseurs au sein d'un même GPU. Elle est dotée d'une large bande passante mémoire (jusqu'à 288 Go/s pour la nouvelle génération Kepler). Cependant, elle possède un temps d'accès plus lent par rapport aux autres mémoires (200 à 600 cycles); \\ \item \textit{Mémoire locale}: de 16 Ko à 512 Ko par processeur. C'est une zone mémoire, accessible en lecture/écriture, dans la mémoire globale. Elle est allouée à un processeur dans le cas où un programme, en cours d'exécution, nécessite plus de registres que ceux disponibles. Bien évidemment, elle possède les mêmes caractéristiques que la mémoire globale; \\ \item \textit{Mémoire constante}: c'est un espace mémoire de 64 Ko qui réside dans la mémoire globale. Il permet de sauvegarder les données dont les valeurs restent constantes au cours de l'exécution d'un programme sur le GPU. De plus, chaque multiprocesseur possède une petite \textit{mémoire cache constante} (environ 8 Ko par multiprocesseur), accessible en lecture seule par tous ses processeurs. Cette mémoire cache constante permet de mettre en cache la mémoire constante, afin d'accélérer les accès mémoires en lecture aux données constantes stockées dans la mémoire constante; \\ \item \textit{Mémoire texture}: n'importe quelle partie de la mémoire globale peut être définie comme une mémoire texture. Elle permet d'améliorer le temps des accès irréguliers à la mémoire globale (voir section~\ref{sec:perf}). Elle peut prendre en charge des tableaux de différents types de données à un, deux ou trois dimensions. Comme pour la mémoire constante, la mémoire texture est mise en cache dans une \textit{mémoire cache texture}, de 6 Ko à 8 Ko par multiprocesseur. Cette mémoire cache texture est accessible en lecture seule par tous les processeurs d'un même multiprocesseur.\\ \end{itemize*} Etant donné que l'espace de la mémoire locale réside dans la mémoire globale, les accès en lecture/écriture à celle-ci ont une latence élevée et une bande passante mémoire faible, par rapport à ceux effectués sur la mémoire partagée. Les nouvelles architectures, ayant une capacité de calcul supérieure ou égale à 2.x (Fermi, Kepler, etc), ont 64 Ko de mémoire par multiprocesseur, configurable en 16 Ko de mémoire partagée et 48 Ko de mémoire cache \textit{L1}, en 48 Ko de mémoire partagée et 16 Ko de mémoire cache \textit{L1} ou 32 Ko de mémoire partagée et 32 Ko de mémoire cache \textit{L1}. De plus, elles possèdent aussi une mémoire cache \textit{L2} de 768 Ko (Fermi) à 1538 Ko (Kepler), partagée entre tous les multiprocesseurs du GPU. Ces deux mémoires caches sont souvent utilisées pour améliorer les performances des accès aux mémoires locale et globale. La seule mémoire GPU accessible par le CPU est la mémoire globale. Tous les échanges de données entre un CPU et son GPU sont effectués via l'interface de communication PCI-Express, de la RAM CPU vers la mémoire globale GPU et vise versa. Ainsi, le CPU peut accéder en lecture/écriture aux mémoires globale, texture et constante. Dans le monde du calcul haute performance, les architectures massivement parallèles des GPUs offrent des performances et des capacités de calcul très intéressantes, pour résoudre de nouveaux problèmes complexes de tailles toujours croissantes. Les deux figures \ref{fig:power} et \ref{fig:bandwidth} montrent, respectivement, la puissance de calcul et la bande passante mémoire théoriques des GPUs Tesla de différentes architectures. La puissance de calcul d'un GPU est représentée par le nombre d'opérations à virgule flottante exécutées par seconde (flops/s). La figure~\ref{fig:power} montre qu'un seul GPU peut fournir une puissance de calcul dépassant les 1 Tflops/s en simple précision ($10^{12}$ flops/s) et les 500 Gflops/s en double précision ($5\times10^{11}$ flops/s). Par ailleurs, une bande passante mémoire exprimée en nombre d'octets par seconde (o/s) désigne le débit de lecture/écriture des données dans la mémoire globale par les processeurs du GPU. La figure~\ref{fig:bandwidth} montre que les bandes passantes mémoires GPU sont très élevées, variant entre 177 Go/s et 288 Go/s, permettant ainsi de diminuer les attentes dues aux accès à la mémoire et augmenter la puissance de calcul. \begin{figure} \centering \includegraphics[width=80mm,keepaspectratio]{Figures/Power/power} \caption{Performance théorique en Gflops/s des GPUs Tesla de différentes architectures.} \label{fig:power} \end{figure} \begin{figure} \centering \includegraphics[width=70mm,keepaspectratio]{Figures/Power/bandwidth} \caption{Bande passante mémoire théorique en Go/s des GPUs Tesla de différentes architectures.} \label{fig:bandwidth} \end{figure} \begin{figure} \centering \includegraphics[width=80mm,keepaspectratio]{Figures/watt} \caption{Rapport performance théorique en double précision et consommation d'énergie en Gflops/Watt.} \label{fig:watt} \end{figure} Un autre paramètre de performance intéressant des GPUs est leur efficacité énergétique. Dans les dernières années, l'architecture des nouveaux produits GPU a été optimisée afin d'augmenter leurs puissances de calcul tout en réduisant leurs consommations d'énergie. La figure~\ref{fig:watt} illustre le rapport entre la puissance de calcul théorique et la consommation énergétique des GPUs de différentes architectures. Ce rapport est exprimé en nombre d'opérations à virgule flottante en double précision exécutées par Watt (flops/Watt). Nous pouvons remarquer que les GPUs des deux premières générations Tesla et Fermi exécutent au maximum 2 Gflops/Watt, alors que ceux des nouvelles générations Kepler et Maxwell, prévues pour 2012 et 2014, pourront exécuter, respectivement, jusqu'à 6 Gflops/Watt et 16 Gflops/Watt en double précision. De quoi intéresser les entreprises et les industries pour réduire les coûts de consommation énergétique de leurs applications. %%-------------------------------------------------------------------------------------------------------%% %%-------------------------------------------------------------------------------------------------------%% \subsection{Programmation multithreadée CUDA} \label{sec:cuda} CUDA est un environnement de programmation des GPUs développé par NVIDIA~\cite{ref19}, dont la première version a été publiée durant l'année 2007. Il est basé sur le langage de programmation C/C++ avec quelques extensions permettant aux GPUs d'exécuter des calculs généraux (applications graphiques et/ou non-graphiques), qui sont habituellement exécutés par les CPUs. Une application écrite en CUDA est un programme hétérogène qui s'exécute sur un processeur (CPU) équipé d'une carte graphique (GPU). En effet, dans un programme CUDA, les codes à exécuter par le CPU sont définis séparément de ceux à exécuter par le GPU. Toutes les opérations à calculs intensifs et faciles à paralléliser sont exécutées par le GPU sous formes de \textit{kernels}. Un kernel est une procédure écrite en CUDA et définie par une entête \verb+__global__+, qui est destinée à être exécutée par le GPU. Par ailleurs, le CPU exécute toutes les opérations séquentielles qui ne peuvent pas être exécutées en parallèle et contrôle l'exécution des kernels sur le GPU ainsi que les communications de données entre la mémoire CPU et la mémoire globale GPU. Au niveau GPU, chaque kernel est exécuté en parallèle par des milliers, voire des millions, de threads. Les threads CUDA d'un même kernel sont organisés en grille de plusieurs blocs de threads, qui sont distribués, plus ou moins équitablement, sur l'ensemble des multiprocesseurs du GPU (voir figure~\ref{fig:block}). En effet, le modèle de programmation CUDA est basé sur un parallélisme de données hiérarchique. Au plus haut niveau, un GPU exécute une grille de blocs de threads en appliquant le modèle parallèle SPMD, où tous les threads exécutent, simultanément, le même code (kernel) mais opérant sur des données différentes. Au niveau intermédiaire, chaque multiprocesseur de GPU exécute un ou plusieurs blocs de threads en appliquant le modèle parallèle SIMD. La position d'un bloc de threads dans la grille est repérée par ses coordonnées à une, deux ou trois dimensions. Au plus bas niveau, chaque c\oe ur d'un multiprocesseur exécute un ou plusieurs threads appartenant au même bloc de threads. A ce niveau, le modèle parallèle SIMT est appliqué de façon à ce que chaque instruction d'un kernel soit exécutée, simultanément, par de multiples threads indépendants (multiples c\oe urs GPU), opérant sur des données différentes. De même que pour les blocs de threads dans une grille, la position d'un thread au sein du bloc, auquel il appartient, est repérée par ses coordonnées à une, deux ou trois dimensions. \begin{figure}[!h] \centering \includegraphics[width=100mm,keepaspectratio]{Figures/block} \caption{Exemple d'exécution des blocs de threads à deux dimensions sur un GPU à 3 multiprocesseurs ayant chacun 8 c\oe urs.} \label{fig:block} \end{figure} \begin{figure}[!h] \centering \includegraphics[width=105mm,keepaspectratio]{Figures/threads} \caption{Exemple d'exécution d'un warp par un multiprocesseur à 8 c\oe urs.} \label{fig:threads} \end{figure} Les threads CUDA peuvent accéder aux différentes mémoires GPU (définies dans la section~\ref{sec:archiGPU}) de manière hiérarchique. Chaque thread a sa propre mémoire locale et ses propres registres. Ensuite, chaque bloc de threads a une mémoire partagée visible par tous ses threads dont la durée de vie des données est la même que celle du bloc de threads. Enfin, tous les threads d'un kernel ont accès à la même mémoire globale et, ainsi, aux mêmes mémoires texture et constante. De plus, dans les nouvelles architectures GPU, tous les threads d'un même bloc partagent une mémoire cache \textit{L1} commune et tous les blocs de threads ont accès à la même mémoire cache \textit{L2}. Au niveau d'un multiprocesseur GPU, les threads d'un même bloc sont exécutés par groupe de 32 threads consécutifs, appelé \textit{warp}. Les threads d'un même warp sont exécutés ensemble, instruction par instruction, jusqu'à la fin du kernel (voir figure~\ref{fig:threads}) et ils sont libres de suivre des chemins d'exécution identiques ou différents, sans aucun point de synchronisation. Au sein d'un même bloc, les threads peuvent coopérer entre eux via la mémoire partagée et synchroniser leurs exécutions en utilisant des barrières de synchronisation (\verb+__syncthreads()+ en CUDA). En revanche, dans la grille de threads d'un kernel, il n'y a aucun moyen de synchronisation entre les différents blocs de threads, si ce n'est qu'ils peuvent seulement lire/écrire dans la même mémoire globale. Le contexte d'exécution (compteurs d'instructions, registres, etc) de chaque warp actif (n'ayant pas encore atteint la fin du kernel) est sauvegardé et maintenu sur le multiprocesseur durant toute la durée de vie du warp. Cela implique que le changement de contexte d'exécution d'un warp à un autre n'a aucune conséquence pénalisant le temps d'exécution d'un kernel. Cependant, cela signifie aussi que tous les warps actifs exécutés par un multiprocesseur partagent les mêmes ressources. Par conséquent, les nombres de threads par bloc et de blocs de threads par grille d'un kernel sont limités par la quantité de ressources disponibles sur un GPU. Un kernel ne peut pas s'exécuter sur un GPU lorsque le nombre de threads par bloc, spécifié par le CPU dans la configuration d'exécution du kernel, est au-dessus du nombre maximum de threads par bloc (512 threads pour Tesla et 1024 threads pour Fermi) ou nécessite plus de registres et/ou d'espace mémoire partagée que disponibles. %%-------------------------------------------------------------------------------------------------------%% %%-------------------------------------------------------------------------------------------------------%% \subsection{Instructions d'optimisation des performances GPU} \label{sec:perf} Pour pouvoir exploiter les performances des GPUs, il est nécessaire, tout d'abord et avant tout, de bien connaître les propriétés de l'architecture matérielle et de l'environnement de programmation des cartes graphiques GPUs utilisées. Par ailleurs, une mise en \oe uvre efficace d'une application sur les GPUs nécessite de bien déterminer les tâches séquentielles et les tâches parallèles de cette application. En effet, toutes les opérations qui sont faciles à exécuter en parallèle doivent être effectuées par le GPU afin d'accélérer l'exécution de l'application. Par contre, toutes les opérations séquentielles et les opérations qui nécessitent des dépendances de données entre threads ou à effectuer des calculs récursifs doivent être exécutées par un seul thread CUDA ou par le CPU, selon la taille du problème à traiter. En fait, l'attente d'un thread pour les résultats de calculs des autres threads affecte considérablement les performances des GPUs. L'efficacité d'un algorithme mis en \oe uvre sur un GPU est étroitement liée à la manière dont les ressources GPU ont été utilisées. Pour optimiser les performances d'un algorithme sur un GPU, il est nécessaire de maximiser l'utilisation des c\oe urs GPU (maximiser le nombre de threads exécutés en parallèle) et d'optimiser l'utilisation des différentes mémoires GPU. \subsubsection{Utilisation des c\oe urs GPU} Comme nous l'avons déjà présenté dans la section~\ref{sec:cuda}, les différents blocs de threads d'un même kernel sont exécutés en parallèle par les différents multiprocesseurs d'un GPU. Afin d'optimiser l'utilisation de ces multiprocesseurs, il convient donc que le nombre de blocs de threads soit un multiple du nombre de multiprocesseurs du GPU utilisé. Ensuite, chaque bloc de threads est partitionné en warps, car un multiprocesseur utilise des warps de 32 threads pour exécuter chaque instruction d'un kernel. Pour maximiser l'utilisation du multiprocesseur, il est nécessaire d'utiliser des multiples de 32 threads pour la taille d'un bloc de threads (32, 64, 128, etc), dans la limite du nombre maximum de threads par bloc. Au niveau d'un multiprocesseur GPU, les différents warps d'un même bloc de threads ne sont pas exécutés en parallèle. Toutefois, lorsque un warp actif doit attendre les données ou le résultat d'une longue opération (par exemple, l'accès à la mémoire globale), il sera mis dans une file d'attente et un autre warp dans la liste des warps prêts (ayant toutes les données nécessaires pour leurs exécutions) sera exécuté. Le nombre de cycles d'horloge nécessaire pour qu'un warp soit prêt à l'exécution est appelé la \textit{latence}. Pour masquer les opérations de grande latence, plus particulièrement les accès à la mémoire globale, un bloc de threads doit avoir plus de 32 threads et donc, au moins deux warps. En outre, les 32 threads d'un même warp exécutent, simultanément, la même instruction d'un kernel (voir section~\ref{sec:cuda}). Donc, l'exécution optimale d'un kernel sur un GPU est assurée lorsque tous les threads d'un même warp suivent le même chemin d'exécution. Dans le cas de divergence d'un warp qui se produit lors des structures conditionnelles (\verb+if(conditions) ... else ...+), le modèle parallèle SIMT force l'évaluation séquentielle des chemins d'exécution des deux branches conditionnelles. En effet, les threads n'entrant pas dans l'une des branches conditionnelles doivent attendre la fin d'exécution des autres threads qui eux, sont entrés dans cette branche. Par conséquence, le temps d'exécution d'une structure conditionnelle est la somme de ceux des chemins d'exécution des différentes branches conditionnelles. \subsubsection{Utilisation des mémoires GPU} Pour maximiser le débit mémoire d'une application mise en \oe uvre sur un GPU, il est nécessaire de réduire les transferts de données entre la mémoire CPU et la mémoire globale GPU, pendant l'exécution de cette application. En raison du surcoût des communications CPU/GPU, il est préférable de regrouper les données de plusieurs petits transferts en un seul grand transfert de données CPU/GPU, que d'effectuer chaque petit transfert séparément. De plus, quand cela est possible, il est intéressant d'utiliser des communications asynchrones entre un CPU et son GPU~\cite{ref19}. Par ailleurs, il est nécessaire aussi de réduire les accès à la mémoire globale et maximiser l'utilisation de la mémoire partagée et des mémoires caches (texture, constante et les caches \textit{L1} et \text{L2} disponibles dans les nouvelles architectures). La mémoire partagée est la mémoire GPU ayant un temps d'accès le plus rapide. Elle est découpée en 16 (Tesla) ou 32 (Fermi et Kepler) modules mémoires de même taille, nommés \textit{banques}, qui peuvent être adressés (lus ou écrits) simultanément par différents threads. Pour optimiser le débit de la mémoire partagée, il faut que $n$ threads d'un même warp puissent accéder, simultanément, à $n$ banques mémoires distinctes ($n$ peut être égal à 16 ou 32, selon l'architecture GPU). Par contre, dans le cas où $m$ différents threads ($m\leq n$) lisent ou écrivent, simultanément, dans la même banque mémoire, leurs accès sont traités en séquentiel, ce qui dégradera les performances. Par ailleurs, la mémoire partagée est souvent exploitée pour le stockage de données réutilisées plusieurs fois au sein d'un même bloc de threads. De plus, pour une utilisation optimale de la mémoire globale, chaque warp doit effectuer, autant que possible, des accès mémoires en lecture/écriture dits \textit{coalescents}. En effet, la mémoire globale est accessible via des transactions mémoires, permettant de lire/écrire des cases voisines alignées sur des segments de 32, 64 ou 128 octets. Selon l'architecture matérielle du GPU, elle peut être adressée par un demi-warp (16 premiers ou 16 derniers threads) dans Tesla ou par un warp dans Fermi et Kepler. Donc, une coalescence absolue est assurée lorsqu'un demi-warp ou un warp accède, simultanément, à 16 ou 32 mots mémoires, respectivement, de même type de données et alignés dans le même segment de la mémoire globale. Le cas contraire se produit lorsqu'un demi-warp (Tesla) ou un warp (Fermi et Kepler) accède à $16$ ou $32$ mots mémoires résidant dans $n$ segments de mémoire différents. Dans ce dernier cas, $n$ transactions mémoires sont nécessaires pour réaliser un seul accès en lecture/écriture pour un demi-warp ou un warp. La figure~\ref{fig:coalescence} montre deux exemples d'accès mémoires: coalescent (figure $(a)$) et non coalescent (figure $(b)$). Pour chaque architecture, nous avons donné le nombre de transactions mémoires nécessaire pour la lecture ou l'écriture de 32 mots mémoires de 4 octets chacun par un warp. Lorsque la coalescence n'est pas assurée, l'utilisation de la mémoire texture est recommandée pour améliorer les temps d'accès en lecture à la mémoire globale. \begin{figure}[!h] \centering \begin{tabular}{c} \includegraphics[width=140mm,keepaspectratio]{Figures/coalescence} \\ \begin{tabular}{|c|c|c|} \hline Architecture & Tesla & Fermi et Kepler \\ \hline Nombre de transactions & 1x de 64 octets à 128 & 1x de 128 octets à 128 \\ mémoires & 1x de 64 octets à 192 & \\ \hline \end{tabular} \\ \\ (a) Accès mémoire coalescent \\ \\ \\ \includegraphics[width=140mm,keepaspectratio]{Figures/coalescence_1} \\ \begin{tabular}{|c|c|c|} \hline Architecture & Tesla & Fermi et Kepler \\ \hline Nombre de transactions & 1x de 128 octets à 128 & 1x de 128 octets à 128 \\ mémoires & 1x de 64 octets à 192 & 1x de 128 octets à 256 \\ & 1x de 32 octets à 256 & \\ \hline \end{tabular} \\ \\ (b) Accès mémoire non coalescent \end{tabular} \caption{Exemples d'accès mémoires coalescent et non coalescent à la mémoire globale par un warp. Un mot mémoire de 4 octets par threads à partir de l'adresse 128.} \label{fig:coalescence} \end{figure} %%-------------------------------------------------------------------------------------------------------%% %%-------------------------------------------------------------------------------------------------------%% \section{Plateformes de calcul parallèle multi-GPUs} \label{sec:multiGPU} %\subsection{Clusters GPU} %\subsection{Grilles de calcul GPU} %More efficient space and energy use — GPU solutions use less space and energy than traditional HPC installations. %GPU solutions either install into standard PCIe slots inside systems or connect externally as a dedicated % server or PCIe chassis using a host interface card. %There is growing interest in building supercomputers that use graphics processors along with CPUs. GPUs are typically faster than traditional CPUs at executing certain tasks, such as those used in scientif%ic and computing applications. Some institutions like the Tokyo Institute of Technology have announced plans to deploy more GPUs in an effort to squeeze more performance out of servers. %%-------------------------------------------------------------------------------------------------------%% %%% Local Variables: %%% mode: latex %%% TeX-master: "these" %%% End: