Figures/CHAPITRE_01.tex~

   1 \chapter{Architectures parall\`eles de GPUs}
   2 \label{chap:archi_parallel}
   3 %Introduction
   4 The high performance computing (HPC) industry’s need for
   5 computation is increasing, as large and complex computational
   6 problems become commonplace across many industry segments.
   7 Traditional CPU technology, however, is no longer capable of scaling
   8 in performance sufficiently to address this demand.
   9 The parallel processing capability of the GPU allows it to divide
  10 complex computing tasks into thousands of smaller tasks that can
  11 be run concurrently. This ability is enabling computational scientists
  12 and researchers to address some of the world’s most challenging
  13 computational problems up to several orders of magnitude faster.
  14
  15 At the end of 2006, NVIDIA GPUs have invaded the field of general purpose computing
  16 with the invention of Compute Unified Device Architecture (CUDA). The performance
  17 and memory bandwidth of the NVIDIA GPUs make them an attractive choice for solving
  18 many scientific computational problems.
  19
  20 In November 2006, NVIDIA introduced CUDATM, a general purpose parallel computing architecture –
  21 with a new parallel programming model and instruction set architecture –
  22 that leverages the parallel compute engine in NVIDIA GPUs to solve many complex computational problems in a more efficient way than on a
  23 CPU.
  24
  25 L'architecture GPU utilisée dans ce document est celle basée sur la plateforme CUDA (Compute Unified Device
  26 Architecture), développée par NVIDIA pour une programmation générique des GPUs afin de bénéficier de leur
  27 capacité de traitement massivement parallèle.
  28
  29 Pourquoi parler des accelerateurs graphiques ?
  30 Ce cours propose de decouvrir et/ou de comprendre les raisons du
  31 succes des accelerateurs de calcul dans le monde du calcul haute
  32 performance (HPC). Depuis plusieurs annees maintenant les
  33 constructeurs esperaient proposer des architectures materielles
  34 simples et puissantes permettant de depasser les limites atteintes
  35 par les CPU. La technologie de GPGPU ( general purpose
  36 processing on graphics processing unit ) s'est imposee au milieu
  37 d'autres technologies existantes pour des raisons d'architecture
  38 materielle mais aussi logicielle. L'objectif double de cours est de
  39 presenter ces elements d'architecture ainsi que l'adequation avec
  40 les problematiques actuelles des thematiques, gourmandes en
  41 puissance de calcul, du monde du HPC.
  42
  43 l'un des plus grands fournisseurs de processeurs graphiques: NVIDIA
  44
  45 définition HPC: high performance computing
  46 %%-------------------------------------------------------------------------------------------------------%%
  47
  48
  49 %%-------------------------------------------------------------------------------------------------------%%
  50 \section{Calcul parallèle}
  51 \label{sec:parallel}
  52 Pour définir le principe d'un calcul parallèle, nous préférons définir, tout d'abord, celui de son opposé:
  53 le calcul séquentiel, pour bien cerner la différence entre ces deux manières de calcul. Un calcul séquentiel
  54 consiste en l'exécution d'un programme, instruction par instruction, par un seul processeur (unité de calcul)
  55 et de façon à ce qu'une seule instruction soit exécutée à la fois. En revanche, un calcul parallèle est défini
  56 comme l'exécution d'un même programme, simultanément, par plusieurs processeurs. Nous avons, en général, deux
  57 façons de réaliser un calcul parallèle. La première consiste en le découpage du programme en plusieurs tâches
  58 de calcul puis, d'exécuter toutes ces tâches en parallèle par différents processeurs. La seconde nécessite
  59 d'abord le partitionnement des données du problème à traiter, de manière à ce que chaque partie de données
  60 soit attribuée à un processeur différent. Ensuite, tous les processeurs exécutent en parallèle les mêmes
  61 instructions du programme mais, opérant sur des données différentes. Cette dernière méthode, appelée la
  62 \textit{parallélisation des données}, est celle retenue dans ce document.
  63
  64 En outre, les calculs parallèles nécessitent aussi une gestion des dépendances de données entre les différents
  65 processeurs. Les calculs locaux de deux processeurs sont dits dépendants lorsque l'exécution de
  66 l'un affecte le résultat de l'autre. Une dépendance de données implique une utilisation de la valeur d'une
  67 même variable par les calculs locaux de deux ou plusieurs processeurs. Les dépendances de données peuvent être gérées
  68 par la synchronisation des lectures/écritures dans une même mémoire (systèmes à mémoires partagées) ou par
  69 la communication de données entre processeurs via des messages (systèmes à mémoires distribuées).
  70
  71 Le calcul parallèle a pour objectif d'exploiter la grande quantité de ressources (puissance de calcul,
  72 espace mémoire, ...) que permettent d'offrir les calculateurs parallèles, pour réduire le temps d'exécution
  73 des applications nécessitant un long traitement et/ou pour pouvoir exécuter celles portant sur des volumes
  74 de données très importants. Ainsi, il nous permet d'aborder de nouveaux problèmes, de plus en plus, complexes
  75 et de tailles toujours croissantes.
  76
  77 \subsection{Classification des architectures parallèles}
  78 Un calculateur parallèle peut être, tout simplement, un processeur multic\oe ures possédant au moins deux unités de calcul physiques
  79 gravées sur la même puce, un supercalculateur qui permet de rassembler les composantes de plusieurs ordinateurs
  80 (processeurs et mémoires) dans une seule machine ou, une plateforme distribuée composée de plusieurs ordinateurs
  81 indépendants, homogènes ou hétérogènes, reliés entre eux par un réseau de communication.
  82
  83 Il existe dans la littérature plusieurs classifications pour les architectures des calculateurs parallèles,
  84 basées sur différents critères de classification~\cite{ref21,ref22,ref23,ref24}. Dans cette section, nous
  85 présentons la classification la plus largement utilisée dans le domaine du calcul parallèle, nommée la
  86 \textit{taxonomie de Flynn}~\cite{ref21}. Elle est basée sur deux critères: le nombre d'instructions et le
  87 nombre de données, qui peuvent être traitées, simultanément, par les différents processeurs du calculateur
  88 parallèle. Les quatre catégories possibles de la taxonomie de Flynn sont les suivantes.
  89
  90 \subsubsection{Instruction unique, donnée unique (SISD)}
  91 La classe SISD (Single Instruction, Single Data) représente l'ensemble des calculateurs séquentiels à
  92 une seule unité de calcul (ou monoprocesseur). Ce sont les calculateurs qui ne sont capables de traiter
  93 qu'une seule instruction sur une seule donnée, par cycle d'horloge. Bien évidemment, cette catégorie n'est
  94 pas une architecture parallèle.
  95
  96 \subsubsection{Instructions multiples, donnée unique (MISD)}
  97 La classe MISD (Multiple Instruction, Single Data) correspond aux calculateurs parallèles pouvant exécuter
  98 plusieurs instructions, simultanément, sur la même donnée. Peu de calculateurs MISD ont existé en pratique,
  99 vu le nombre réduit des applications qui peuvent être mises en \oe uvre sur ce type d'architecture. Un
 100 exemple de calculateur parallèle expérimental MISD a été développé à l'université de Carnegie Mellon~\cite{ref25}.
 101
 102 \subsubsection{Instruction unique, données multiples (SIMD)}
 103 La classe SIMD (Single Instruction, Multiple Data) correspond aux processeurs vectoriels et, plus généralement,
 104 aux calculateurs composés d'un grand nombre d'unités de calcul. Chaque processeur d'un calculateur SIMD exécute
 105 la même instruction à chaque cycle d'horloge, mais opérant sur des données différentes. Cette architecture
 106 parallèle est bien adaptée aux traitement des problèmes à structure régulière, où la même instruction est
 107 appliquée à un ensemble de données (exécutions des opérations sur des vecteurs ou des tableaux).
 108
 109 \subsubsection{Instructions multiples, données multiples (MIMD)}
 110 La classe MIMD (Multiple Instruction, Multiple Data) représente la classe la plus générale dans cette
 111 classification. Les calculateurs parallèles MIMD possèdent plusieurs processeurs interconnectés entre
 112 eux, tels que chaque processeur est capable de suivre son propre chemin d'exécution. En effet, à chaque
 113 cycle d'horloge, les processeurs peuvent exécuter des instructions différentes sur des données différentes.
 114
 115 \subsection{Mémoires des architectures parallèles}
 116 Nous pouvons distinguer, en général, deux modèles de gestion de la mémoire des calculateurs parallèles.
 117
 118 \subsubsection{Mémoire partagée}
 119 Dans ce type d'architecture,
 120
 121 \subsubsection{Mémoire distribuée}
 122 %%-------------------------------------------------------------------------------------------------------%%
 123
 124
 125 %%-------------------------------------------------------------------------------------------------------%%
 126 \section{Unité de traitement graphique GPU}
 127 \label{sec:GPU}
 128 L'architecture et l'environnement de programmation des GPUs utilisés dans ce document sont ceux basés sur la
 129 plateforme CUDA (Compute Unified Device Architecture) développée par NVIDIA~\cite{ref19}.
 130
 131 \subsection{Architecture matérielle GPU}
 132 \label{sec:archiGPU}
 133 Les processeurs graphiques GPUs sont initialement conçus pour le traitement des applications graphiques et
 134 de la visualisation 3D. Nous pouvons citer, par exemple, les produits \textit{GeForce} et \textit{Quadro},
 135 deux gammes de GPUs proposées par NVIDIA, qui sont destinés, respectivement, au graphisme grand public et à
 136 la visualisation professionnelle. Depuis quelques années, les GPUs sont devenus des outils très attrayants
 137 pour le calcul haute performance (HPC). La gamme de produits \textit{Tesla} a été conçue par NVIDIA pour
 138 offrir des capacités de calcul parallèle élevées et assister les processeurs dans les calculs intensifs des
 139 applications scientifiques et/ou industrielles. La figure~\ref{fig:archi} montre les différentes architectures
 140 matérielles GPU développées par NVIDIA.
 141 \begin{figure}[!h]
 142   \centering
 143   \includegraphics[width=85mm,keepaspectratio]{Figures/archi}
 144 \caption{Historique des architectures matérielles GPU.}
 145 \label{fig:archi}
 146 \end{figure}
 147
 148 Un GPU est un processeur graphique relié à un processeur traditionnel (CPU) via un PCI-Express
 149 (voir figure~\ref{fig:gpu-cpu}). Il est souvent considéré comme un accélérateur des tâches
 150 parallèles et des opérations arithmétiques intensives d'une application exécutée sur un CPU.
 151 Il puise sa puissance de calcul de son architecture matérielle et logicielle massivement parallèle.
 152 En effet, à la différence d'une architecture CPU, un GPU est composé de centaines (voire de
 153 milliers) de processeurs (SP), appelés communément c\oe urs, organisés en plusieurs blocs de
 154 processeurs appelés multiprocesseurs (SM ou SMX). La figure~\ref{fig:compar} montre une comparaison
 155 entre l'architecture matérielle d'un CPU et celle d'un GPU Fermi. Les processeurs d'un GPU sont
 156 regroupés par 8 (Tesla), 32 (Fermi) ou 192 (Kepler) dans un multiprocesseur, selon le type de
 157 son architecture matérielle. De la même manière, les multiprocesseurs sont eux-mêmes regroupés
 158 par 2 (G80) ou 3 (GT200) dans un TPC (Texture Processing Cluster) pour l'architecture Tesla et
 159 par 4 (Fermi) ou 2 (Kepler) dans un GPC (Graphics Processing Cluster) pour les nouvelles architectures.
 160 \begin{figure}[!h]
 161   \centering
 162    \begin{tabular}{ccc}
 163       \includegraphics[width=45mm,keepaspectratio]{Figures/fig} &
 164       \includegraphics[width=50mm,keepaspectratio]{Figures/fig1} &
 165       \includegraphics[width=60mm,keepaspectratio]{Figures/schema1} \\
 166       (a) Carte graphique GPU & (b) Un GPU relié à un CPU & (c) Un schéma de GPU relié à un CPU\\
 167       \end{tabular}
 168   \caption{Un exemple de CPU équipé d'un GPU}
 169   \label{fig:gpu-cpu}
 170 \end{figure}
 171
 172 \begin{figure}[!h]
 173   \centering
 174   \begin{tabular}{cc}
 175     \includegraphics[width=50mm,keepaspectratio]{Figures/CPU_scheme} &
 176     \includegraphics[width=50mm,keepaspectratio]{Figures/GPU_scheme} \\
 177     (a) Un CPU à 8 c\oe urs & (b) Un GPU Fermi à 512 c\oe urs
 178    \end{tabular}
 179   \caption{Comparaison du nombre de c\oe urs dans un CPU et dans un GPU.}
 180   \label{fig:compar}
 181 \end{figure}
 182
 183 En plus de la hiérarchie de processeurs, un GPU est doté d'une hiérarchie de mémoires de différentes
 184 tailles et de différentes bandes passantes mémoires. Nous distinguons, au total, six mémoires différentes
 185 (voir figure~\ref{fig:memoires}):
 186
 187 \begin{figure}[!h]
 188   \centering
 189   \begin{tabular}{cc}
 190      \includegraphics[width=82mm,keepaspectratio]{Figures/memoires} &
 191      \includegraphics[width=82mm,keepaspectratio]{Figures/memoiresFermi} \\
 192      (a) Architecture Tesla & (b) Architecture Fermi ou Kepler \\
 193   \end{tabular}
 194   \caption{Hiérarchie de mémoires GPU.}
 195   \label{fig:memoires}
 196 \end{figure}
 197
 198 \begin{itemize*}
 199   \item \textit{Registres}: chaque multiprocesseur a 8K à 65K de registres à 32-bit, répartis entre tous ses
 200 processeurs. Ce sont des mémoires rapides, accessibles en lecture/écriture et avec une faible latence (environ
 201 1 cycle);
 202 \\
 203   \item \textit{Mémoire partagée}: de 16 Ko à 48 Ko de mémoire par multiprocesseur. C'est une petite mémoire
 204 extrêmement rapide. Elle est dotée d'une large bande passante mémoire (plus d'un To/s) et d'une faible latence
 205 (environ 1 à 2 cycles). Elle est accessible en lecture/écriture par tous les processeurs du même multiprocesseur;
 206 \\
 207   \item \textit{Mémoire globale}: chaque GPU est équipé de sa propre RAM (GDDR3 ou GDDR5) de 1 Go à 6 Go. C'est
 208 une mémoire accessible en lecture/écriture et partagée entre tous les multiprocesseurs au sein d'un même GPU.
 209 Elle est dotée d'une large bande passante mémoire (jusqu'à 288 Go/s pour la nouvelle génération Kepler). Cependant,
 210 elle possède un temps d'accès plus lent par rapport aux autres mémoires (200 à 600 cycles);
 211 \\
 212   \item \textit{Mémoire locale}: de 16 Ko à 512 Ko par processeur. C'est une zone mémoire, accessible en lecture/écriture,
 213 dans la mémoire globale. Elle est allouée à un processeur dans le cas où un programme, en cours d'exécution,
 214 nécessite plus de registres que ceux disponibles. Bien évidemment, elle possède les mêmes caractéristiques que la mémoire
 215 globale;
 216 \\
 217   \item \textit{Mémoire constante}: c'est un espace mémoire de 64 Ko qui réside dans la mémoire globale. Il
 218 permet de sauvegarder les données dont les valeurs restent constantes au cours de l'exécution d'un programme
 219 sur le GPU. De plus, chaque multiprocesseur possède une petite \textit{mémoire cache constante} (environ 8 Ko
 220 par multiprocesseur), accessible en lecture seule par tous ses processeurs. Cette mémoire cache constante
 221 permet de mettre en cache la mémoire constante, afin d'accélérer les accès mémoires en lecture aux données
 222 constantes stockées dans la mémoire constante;
 223 \\
 224   \item \textit{Mémoire texture}: n'importe quelle partie de la mémoire globale peut être définie comme une
 225 mémoire texture. Elle permet d'améliorer le temps des accès irréguliers à la mémoire globale (voir section~\ref{sec:perf}).
 226 Elle peut prendre en charge des tableaux de différents types de données à un, deux ou trois dimensions. Comme
 227 pour la mémoire constante, la mémoire texture est mise en cache dans une \textit{mémoire cache texture}, de
 228 6 Ko à 8 Ko par multiprocesseur. Cette mémoire cache texture est accessible en lecture seule par tous les processeurs
 229 d'un même multiprocesseur.\\
 230 \end{itemize*}
 231
 232 Etant donné que l'espace de la mémoire locale réside dans la mémoire globale, les accès en lecture/écriture
 233 à celle-ci ont une latence élevée et une bande passante mémoire faible, par rapport à ceux effectués sur la
 234 mémoire partagée. Les nouvelles architectures, ayant une capacité de calcul supérieure ou égale à 2.x (Fermi,
 235 Kepler, etc), ont 64 Ko de mémoire par multiprocesseur, configurable en 16 Ko de mémoire partagée et 48 Ko de
 236 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
 237 partagée et 32 Ko de mémoire cache \textit{L1}. De plus, elles possèdent aussi une mémoire cache \textit{L2}
 238 de 768 Ko (Fermi) à 1538 Ko (Kepler), partagée entre tous les multiprocesseurs du GPU. Ces deux mémoires caches
 239 sont souvent utilisées pour améliorer les performances des accès aux mémoires locale et globale. La seule mémoire
 240 GPU accessible par le CPU est la mémoire globale. Tous les échanges de données entre un CPU et son GPU sont
 241 effectués via l'interface de communication PCI-Express, de la RAM CPU vers la mémoire globale GPU et vise versa.
 242 Ainsi, le CPU peut accéder en lecture/écriture aux mémoires globale, texture et constante.
 243
 244 Dans le monde du calcul haute performance, les architectures massivement parallèles des GPUs offrent des
 245 performances et des capacités de calcul très intéressantes, pour résoudre de nouveaux problèmes complexes
 246 de tailles toujours croissantes. Les deux figures \ref{fig:power} et \ref{fig:bandwidth} montrent, respectivement,
 247 la puissance de calcul et la bande passante mémoire théoriques des GPUs Tesla de différentes architectures.
 248 La puissance de calcul d'un GPU est représentée par le nombre d'opérations à virgule flottante exécutées
 249 par seconde (flops/s). La figure~\ref{fig:power} montre qu'un seul GPU peut fournir une puissance de calcul
 250 dépassant les 1 Tflops/s en simple précision ($10^{12}$ flops/s) et les 500 Gflops/s en double précision
 251 ($5\times10^{11}$ flops/s). Par ailleurs, une bande passante mémoire exprimée en nombre d'octets par seconde
 252 (o/s) désigne le débit de lecture/écriture des données dans la mémoire globale par les processeurs du GPU.
 253 La figure~\ref{fig:bandwidth} montre que les bandes passantes mémoires GPU sont très élevées, variant entre
 254 177 Go/s et 288 Go/s, permettant ainsi de diminuer les attentes dues aux accès à la mémoire et augmenter
 255 la puissance de calcul.
 256
 257 \begin{figure}
 258  \centering
 259   \includegraphics[width=80mm,keepaspectratio]{Figures/Power/power}
 260 \caption{Performance théorique en Gflops/s des GPUs Tesla de différentes architectures.}
 261 \label{fig:power}
 262 \end{figure}
 263
 264 \begin{figure}
 265  \centering
 266   \includegraphics[width=70mm,keepaspectratio]{Figures/Power/bandwidth}
 267 \caption{Bande passante mémoire théorique en Go/s des GPUs Tesla de différentes architectures.}
 268 \label{fig:bandwidth}
 269 \end{figure}
 270
 271 \begin{figure}
 272   \centering
 273   \includegraphics[width=80mm,keepaspectratio]{Figures/watt}
 274 \caption{Rapport performance théorique en double précision et consommation d'énergie en Gflops/Watt.}
 275 \label{fig:watt}
 276 \end{figure}
 277
 278 Un autre paramètre de performance intéressant des GPUs est leur efficacité énergétique. Dans les
 279 dernières années, l'architecture des nouveaux produits GPU a été optimisée afin d'augmenter leurs
 280 puissances de calcul tout en réduisant leurs consommations d'énergie. La figure~\ref{fig:watt}
 281 illustre le rapport entre la puissance de calcul théorique et la consommation énergétique des GPUs
 282 de différentes architectures. Ce rapport est exprimé en nombre d'opérations à virgule flottante en
 283 double précision exécutées par Watt (flops/Watt). Nous pouvons remarquer que les GPUs des deux premières
 284 générations Tesla et Fermi exécutent au maximum 2 Gflops/Watt, alors que ceux des nouvelles générations
 285 Kepler et Maxwell, prévues pour 2012 et 2014, pourront exécuter, respectivement, jusqu'à 6 Gflops/Watt
 286 et 16 Gflops/Watt en double précision. De quoi intéresser les entreprises et les industries pour
 287 réduire les coûts de consommation énergétique de leurs applications.
 288 %%-------------------------------------------------------------------------------------------------------%%
 289
 290
 291 %%-------------------------------------------------------------------------------------------------------%%
 292 \subsection{Programmation multithreadée CUDA}
 293 \label{sec:cuda}
 294 CUDA est un environnement de programmation des GPUs développé par
 295 NVIDIA~\cite{ref19}, dont la première version a été publiée durant l'année 2007. Il est basé sur le langage
 296 de programmation C/C++ avec quelques extensions permettant aux GPUs d'exécuter des calculs généraux (applications
 297 graphiques et/ou non-graphiques), qui sont habituellement exécutés par les CPUs. Une application écrite en
 298 CUDA est un programme hétérogène qui s'exécute sur un processeur (CPU) équipé d'une carte graphique (GPU).
 299 En effet, dans un programme CUDA, les codes à exécuter par le CPU sont définis séparément de ceux à exécuter
 300 par le GPU. Toutes les opérations à calculs intensifs et faciles à paralléliser sont exécutées par le GPU
 301 sous formes de \textit{kernels}. Un kernel est une procédure écrite en CUDA et définie par une entête
 302 \verb+__global__+, qui est destinée à être exécutée par le GPU. Par ailleurs, le CPU exécute toutes les
 303 opérations séquentielles qui ne peuvent pas être exécutées en parallèle et contrôle l'exécution des kernels
 304 sur le GPU ainsi que les communications de données entre la mémoire CPU et la mémoire globale GPU.
 305
 306 Au niveau GPU, chaque kernel est exécuté en parallèle par des milliers, voire des millions, de threads.
 307 Les threads CUDA d'un même kernel sont organisés en grille de plusieurs blocs de threads, qui sont
 308 distribués, plus ou moins équitablement, sur l'ensemble des multiprocesseurs du GPU (voir figure~\ref{fig:block}).
 309 En effet, le modèle de programmation CUDA est basé sur un parallélisme de données hiérarchique. Au plus
 310 haut niveau, un GPU exécute une grille de blocs de threads en appliquant le modèle parallèle SPMD, où
 311 tous les threads exécutent, simultanément, le même code (kernel) mais opérant sur des données différentes.
 312 Au niveau intermédiaire, chaque multiprocesseur de GPU exécute un ou plusieurs blocs de threads en appliquant
 313 le modèle parallèle SIMD. La position d'un bloc de threads dans la grille est repérée par ses coordonnées
 314 à une, deux ou trois dimensions. Au plus bas niveau, chaque c\oe ur d'un multiprocesseur exécute un ou
 315 plusieurs threads appartenant au même bloc de threads. A ce niveau, le modèle parallèle SIMT est appliqué
 316 de façon à ce que chaque instruction d'un kernel soit exécutée, simultanément, par de multiples threads indépendants
 317 (multiples c\oe urs GPU), opérant sur des données différentes. De même que pour les blocs de threads dans
 318 une grille, la position d'un thread au sein du bloc, auquel il appartient, est repérée par ses coordonnées
 319 à une, deux ou trois dimensions.
 320 \begin{figure}[!h]
 321  \centering
 322   \includegraphics[width=100mm,keepaspectratio]{Figures/block}
 323 \caption{Exemple d'exécution des blocs de threads à deux dimensions sur un GPU à 3 multiprocesseurs ayant chacun 8 c\oe urs.}
 324 \label{fig:block}
 325 \end{figure}
 326
 327 \begin{figure}[!h]
 328  \centering
 329   \includegraphics[width=105mm,keepaspectratio]{Figures/threads}
 330 \caption{Exemple d'exécution d'un warp par un multiprocesseur à 8 c\oe urs.}
 331 \label{fig:threads}
 332 \end{figure}
 333
 334 Les threads CUDA peuvent accéder aux différentes mémoires GPU (définies dans la section~\ref{sec:archiGPU})
 335 de manière hiérarchique. Chaque thread a sa propre mémoire locale et ses propres registres. Ensuite, chaque
 336 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
 337 même que celle du bloc de threads. Enfin, tous les threads d'un kernel ont accès à la même mémoire globale
 338 et, ainsi, aux mêmes mémoires texture et constante. De plus, dans les nouvelles architectures GPU, tous les
 339 threads d'un même bloc partagent une mémoire cache \textit{L1} commune et tous les blocs de threads ont accès
 340 à la même mémoire cache \textit{L2}.
 341
 342 Au niveau d'un multiprocesseur GPU, les threads d'un même bloc sont exécutés par groupe de 32 threads consécutifs,
 343 appelé \textit{warp}. Les threads d'un même warp sont exécutés ensemble, instruction par instruction, jusqu'à la fin
 344 du kernel (voir figure~\ref{fig:threads}) et ils sont libres de suivre des chemins d'exécution identiques ou différents,
 345 sans aucun point de synchronisation. Au sein d'un même bloc, les threads peuvent coopérer entre eux via la mémoire
 346 partagée et synchroniser leurs exécutions en utilisant des barrières de synchronisation (\verb+__syncthreads()+
 347 en CUDA). En revanche, dans la grille de threads d'un kernel, il n'y a aucun moyen de synchronisation entre les
 348 différents blocs de threads, si ce n'est qu'ils peuvent seulement lire/écrire dans la même mémoire globale.
 349
 350 Le contexte d'exécution (compteurs d'instructions, registres, etc) de chaque warp actif (n'ayant pas encore
 351 atteint la fin du kernel) est sauvegardé et maintenu sur le multiprocesseur durant toute la durée de vie du
 352 warp. Cela implique que le changement de contexte d'exécution d'un warp à un autre n'a aucune conséquence
 353 pénalisant le temps d'exécution d'un kernel. Cependant, cela signifie aussi que tous les warps actifs exécutés
 354 par un multiprocesseur partagent les mêmes ressources. Par conséquent, les nombres de threads par bloc et de
 355 blocs de threads par grille d'un kernel sont limités par la quantité de ressources disponibles sur un GPU.
 356 Un kernel ne peut pas s'exécuter sur un GPU lorsque le nombre de threads par bloc, spécifié par le CPU dans
 357 la configuration d'exécution du kernel, est au-dessus du nombre maximum de threads par bloc (512 threads pour
 358 Tesla et 1024 threads pour Fermi) ou nécessite plus de registres et/ou d'espace mémoire partagée que disponibles.
 359 %%-------------------------------------------------------------------------------------------------------%%
 360
 361
 362 %%-------------------------------------------------------------------------------------------------------%%
 363 \subsection{Instructions d'optimisation des performances GPU}
 364 \label{sec:perf}
 365 Pour pouvoir exploiter les performances des GPUs, il est nécessaire, tout d'abord et avant tout, de bien
 366 connaître les propriétés de l'architecture matérielle et de l'environnement de programmation des cartes
 367 graphiques GPUs utilisées. Par ailleurs, une mise en \oe uvre efficace d'une application sur les GPUs
 368 nécessite de bien déterminer les tâches séquentielles et les tâches parallèles de cette application.
 369 En effet, toutes les opérations qui sont faciles à exécuter en parallèle doivent être effectuées par
 370 le GPU afin d'accélérer l'exécution de l'application. Par contre, toutes les opérations séquentielles
 371 et les opérations qui nécessitent des dépendances de données entre threads ou à effectuer des calculs
 372 récursifs doivent être exécutées par un seul thread CUDA ou par le CPU, selon la taille du problème à
 373 traiter. En fait, l'attente d'un thread pour les résultats de calculs des autres threads affecte
 374 considérablement les performances des GPUs.
 375
 376 L'efficacité d'un algorithme mis en \oe uvre sur un GPU est étroitement liée à la manière dont les ressources
 377 GPU ont été utilisées. Pour optimiser les performances d'un algorithme sur un GPU, il est nécessaire de
 378 maximiser l'utilisation des c\oe urs GPU (maximiser le nombre de threads exécutés en parallèle) et d'optimiser
 379 l'utilisation des différentes mémoires GPU.
 380
 381 \subsubsection{Utilisation des c\oe urs GPU}
 382 Comme nous l'avons déjà présenté dans la section~\ref{sec:cuda}, les différents blocs de threads d'un même
 383 kernel sont exécutés en parallèle par les différents multiprocesseurs d'un GPU. Afin d'optimiser l'utilisation
 384 de ces multiprocesseurs, il convient donc que le nombre de blocs de threads soit un multiple du
 385 nombre de multiprocesseurs du GPU utilisé. Ensuite, chaque bloc de threads est partitionné en warps, car un
 386 multiprocesseur utilise des warps de 32 threads pour exécuter chaque instruction d'un kernel. Pour maximiser
 387 l'utilisation du multiprocesseur, il est nécessaire d'utiliser des multiples de 32 threads pour la taille d'un
 388 bloc de threads (32, 64, 128, etc), dans la limite du nombre maximum de threads par bloc.
 389
 390 Au niveau d'un multiprocesseur GPU, les différents warps d'un même bloc de threads ne sont pas exécutés en
 391 parallèle. Toutefois, lorsque un warp actif doit attendre les données ou le résultat d'une longue opération
 392 (par exemple, l'accès à la mémoire globale), il sera mis dans une file d'attente et un autre warp dans la
 393 liste des warps prêts (ayant toutes les données nécessaires pour leurs exécutions) sera exécuté.
 394 Le nombre de cycles d'horloge nécessaire pour qu'un warp soit prêt à l'exécution est appelé la \textit{latence}.
 395 Pour masquer les opérations de grande latence, plus particulièrement les accès à la mémoire globale, un bloc
 396 de threads doit avoir plus de 32 threads et donc, au moins deux warps.
 397
 398 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}).
 399 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
 400 même chemin d'exécution. Dans le cas de divergence d'un warp qui se produit lors des structures conditionnelles
 401 (\verb+if(conditions) ... else ...+), le modèle parallèle SIMT force l'évaluation séquentielle des chemins
 402 d'exécution des deux branches conditionnelles. En effet, les threads n'entrant pas dans l'une des branches
 403 conditionnelles doivent attendre la fin d'exécution des autres threads qui eux, sont entrés dans cette branche.
 404 Par conséquence, le temps d'exécution d'une structure conditionnelle est la somme de ceux des chemins d'exécution
 405 des différentes branches conditionnelles.
 406
 407 \subsubsection{Utilisation des mémoires GPU}
 408 Pour maximiser le débit mémoire d'une application mise en \oe uvre sur un GPU, il est nécessaire de réduire
 409 les transferts de données entre la mémoire CPU et la mémoire globale GPU, pendant l'exécution de cette
 410 application. En raison du surcoût des communications CPU/GPU, il est préférable de regrouper les données de
 411 plusieurs petits transferts en un seul grand transfert de données CPU/GPU, que d'effectuer chaque petit
 412 transfert séparément. De plus, quand cela est possible, il est intéressant d'utiliser des communications
 413 asynchrones entre un CPU et son GPU~\cite{ref19}.
 414
 415 Par ailleurs, il est nécessaire aussi de réduire les accès à la mémoire globale et maximiser l'utilisation
 416 de la mémoire partagée et des mémoires caches (texture, constante et les caches \textit{L1} et \text{L2}
 417 disponibles dans les nouvelles architectures). La mémoire partagée est la mémoire GPU ayant un temps d'accès
 418 le plus rapide. Elle est découpée en 16 (Tesla) ou 32 (Fermi et Kepler) modules mémoires de même taille,
 419 nommés \textit{banques}, qui peuvent être adressés (lus ou écrits) simultanément par différents threads.
 420 Pour optimiser le débit de la mémoire partagée, il faut que $n$ threads d'un même warp puissent accéder,
 421 simultanément, à $n$ banques mémoires distinctes ($n$ peut être égal à 16 ou 32, selon l'architecture GPU).
 422 Par contre, dans le cas où $m$ différents threads ($m\leq n$) lisent ou écrivent, simultanément, dans la
 423 même banque mémoire, leurs accès sont traités en séquentiel, ce qui dégradera les performances. Par ailleurs,
 424 la mémoire partagée est souvent exploitée pour le stockage de données réutilisées plusieurs fois au sein
 425 d'un même bloc de threads.
 426
 427 De plus, pour une utilisation optimale de la mémoire globale, chaque warp doit effectuer, autant que possible,
 428 des accès mémoires en lecture/écriture dits \textit{coalescents}. En effet, la mémoire globale est accessible
 429 via des transactions mémoires, permettant de lire/écrire des cases voisines alignées sur des segments de 32,
 430 64 ou 128 octets. Selon l'architecture matérielle du GPU, elle peut être adressée par un demi-warp (16 premiers
 431 ou 16 derniers threads) dans Tesla ou par un warp dans Fermi et Kepler. Donc, une coalescence absolue est
 432 assurée lorsqu'un demi-warp ou un warp accède, simultanément, à 16 ou 32 mots mémoires, respectivement,
 433 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
 434 lorsqu'un demi-warp (Tesla) ou un warp (Fermi et Kepler) accède à $16$ ou $32$ mots mémoires résidant
 435 dans $n$ segments de mémoire différents. Dans ce dernier cas, $n$ transactions mémoires sont nécessaires
 436 pour réaliser un seul accès en lecture/écriture pour un demi-warp ou un warp. La figure~\ref{fig:coalescence}
 437 montre deux exemples d'accès mémoires: coalescent (figure $(a)$) et non coalescent (figure $(b)$). Pour chaque
 438 architecture, nous avons donné le nombre de transactions mémoires nécessaire pour la lecture ou l'écriture
 439 de 32 mots mémoires de 4 octets chacun par un warp. Lorsque la coalescence n'est pas assurée, l'utilisation de
 440 la mémoire texture est recommandée pour améliorer les temps d'accès en lecture à la mémoire globale.
 441
 442 \begin{figure}[!h]
 443   \centering
 444   \begin{tabular}{c}
 445      \includegraphics[width=140mm,keepaspectratio]{Figures/coalescence} \\
 446      \begin{tabular}{|c|c|c|}
 447       \hline
 448       Architecture & Tesla & Fermi et Kepler \\
 449       \hline
 450       Nombre de transactions  & 1x de 64 octets à 128 & 1x de 128 octets à 128 \\
 451       mémoires                & 1x de 64 octets à 192 &                        \\
 452       \hline
 453      \end{tabular}
 454      \\ \\
 455      (a) Accès mémoire coalescent \\ \\ \\
 456
 457      \includegraphics[width=140mm,keepaspectratio]{Figures/coalescence_1} \\
 458      \begin{tabular}{|c|c|c|}
 459       \hline
 460       Architecture & Tesla & Fermi et Kepler \\
 461       \hline
 462       Nombre de transactions  & 1x de 128 octets à 128 & 1x de 128 octets à 128 \\
 463       mémoires                & 1x de 64 octets à 192  & 1x de 128 octets à 256 \\
 464                               & 1x de 32 octets à 256  &                        \\
 465       \hline
 466      \end{tabular}
 467      \\ \\
 468      (b) Accès mémoire non coalescent
 469   \end{tabular}
 470   \caption{Exemples d'accès mémoires coalescent et non coalescent à la mémoire globale par un warp. Un mot
 471     mémoire de 4 octets par threads à partir de l'adresse 128.}
 472   \label{fig:coalescence}
 473 \end{figure}
 474
 475 %%-------------------------------------------------------------------------------------------------------%%
 476
 477
 478 %%-------------------------------------------------------------------------------------------------------%%
 479 \section{Plateformes de calcul parallèle multi-GPUs}
 480 \label{sec:multiGPU}
 481
 482 %\subsection{Clusters GPU}
 483
 484 %\subsection{Grilles de calcul GPU}
 485
 486 %More efficient space and energy use — GPU solutions use less space and energy than traditional HPC installations.
 487 %GPU solutions either install into standard PCIe slots inside systems or connect externally as a dedicated
 488 % server or PCIe chassis using a host interface card.
 489
 490 %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.
 491 %%-------------------------------------------------------------------------------------------------------%%
 492
 493 %%% Local Variables:
 494 %%% mode: latex
 495 %%% TeX-master: "these"
 496 %%% End: