developpement:activites:qualification:30ans1code

Différences

Ci-dessous, les différences entre deux révisions de la page.

--- developpement:activites:qualification:30ans1code [2019/12/17 11:27]
equemene
+++ developpement:activites:qualification:30ans1code [2019/12/17 15:34] (Version actuelle)
equemene
@@ Ligne 1: / Ligne 1: @@
 ====== Simuler un trou noir : du processeur à la carte graphique ======
-<note important>En construction...</note>
 **Simuler l'apparence d'un trou noir et son écharpe de plasma,\\ de la parallélisation à la GéPUfication d'un code,\\ 40 ans d'histoire scientifique en 30 ans d'évolution technologique.**
@@ Ligne 176: / Ligne 174: @@
 {{:developpement:activites:qualification:k6-2_hammsqueeze.png?500|}}
-Si nous enlevons l'augmentation de fréquence (presque doublée), le changement d'architecture du Amd5x86 au K6 apporte un gain de 75% en BB et 2 en Mono si nous restons en 1998 (avec la distribution équipée d'un compilateur GCC version 2.7). Par contre, en 2011 avec un compilateur GCC 4.4, ce gain passe à presque 10 en BB et presque 5 en Mono. Ainsi, les microarchitectures anciennes (et là probablement la vectorisation) sont bien mieux supportées dans les compilateurs récents.
+Si nous enlevons l'augmentation de fréquence (presque doublée), le changement d'architecture du Amd5x86 au K6 apporte un gain de 75% en BB et 2 en Mono, si nous restons en 1998 (avec la distribution équipée d'un compilateur GCC version 2.7). Par contre, en 2011 avec un compilateur GCC 4.4, ce gain passe à presque 10 en BB et presque 5 en Mono. Ainsi, les microarchitectures anciennes (et là probablement la vectorisation) sont bien mieux supportées dans les compilateurs récents.
-Au final, le RISC86 et la vectorisation du K6-2 auront offerts un gain de 2400 en BB et 2200 en Mono. Avec l'augmentation de fréquence, c'est un gain entre 12500 et 14000. C'est plus d'un ordre de grandeur par rapport à ce qu'offrait déjà l'intégration du FPU dans le processeur.
+Au final, le RISC86 et la vectorisation du K6-2 auront offert un gain de 2400 en BB et 2200 en Mono. Avec l'augmentation de fréquence, c'est un gain entre 12500 et 14000. C'est plus d'un ordre de grandeur par rapport à ce qu'offrait déjà l'intégration du FPU dans le processeur.
 === Avant la multiplication des coeurs, le passage de l'an 2000 ===
@@ Ligne 206: / Ligne 204: @@
 {{:developpement:activites:qualification:quadcores.png?500|}}
-Nous constatons non sans une certaine surprise que le processeur le plus récent avec la fréquence la plus élevée (le Skylake Gold 5122 à 3.6 GHz) est bien inférieur à son prédecesseur (le Broadwell E5-2637v4 à 3.5 GHz). Il faut aussi noter que, finalement, entre le Harepertown de 2008 et le Skylake de 2018, sur ce code séquentiel, c'est seulement un gain d'un facteur 3, bien en dessous de ce que 10 années précédentes avaient offert...
+Nous constatons non sans une certaine surprise que le processeur le plus récent avec la fréquence la plus élevée (le Skylake Gold 5122 à 3.6 GHz) est bien inférieur à son prédécesseur (le Broadwell E5-2637v4 à 3.5 GHz). Il faut aussi noter que, finalement, entre le Harepertown de 2008 et le Skylake de 2018, sur ce code séquentiel, c'est seulement un gain d'un facteur 3, bien en dessous de ce que 10 années précédentes avaient offert...
 {{:developpement:activites:qualification:quadcores_log.png?500|}}
@@ Ligne 212: / Ligne 210: @@
 === Avec les processeurs les plus "gros" ou les plus récents ===
-A ces 4 systèmes équipés de 2 processeurs disposant de 4 coeurs, nous ajoutons 4 autres processeurs : un système disposant de 28 coeurs physiques (2 Broadwell E5-2680 à 2.4 GHz), un système disposant de 20 coeurs physiques (2 Skylake Silver à 2.2 GHz), un AMD Threadripper 1950X à 3.6 GHz disposant de 16 coeurs physique et un Skylake W-2145 à 3.6 GHz disposant de 8 coeurs. Ces quatre systèmes illustrent des configurations avec beaucoup de coeurs ou les dernières générations de processeurs disponibles au CBP.
+A ces 4 systèmes équipés de 2 processeurs disposant de 4 coeurs, nous ajoutons 4 autres processeurs : un système disposant de 28 coeurs physiques (2 Broadwell E5-2680 à 2.4 GHz), un système disposant de 20 coeurs physiques (2 Skylake Silver à 2.2 GHz), un AMD Threadripper 1950X à 3.6 GHz disposant de 16 coeurs physiques et un Skylake W-2145 à 3.6 GHz disposant de 8 coeurs. Ces quatre systèmes illustrent des configurations avec beaucoup de coeurs ou les dernières générations de processeurs disponibles au CBP.
  {{:developpement:activites:qualification:serialallprocessors.png?500|}}
@@ Ligne 218: / Ligne 216: @@
 Pas de grande surprise sinon la victoire d'une courte tête du Threadripper 1950X, devant le processeur de station de travail W-2145. Ainsi, ce n'est pas parce que c'est un processeur de serveur que c'est le plus performant, ou que c'est le plus récent à la fréquence la plus élevée.
-Au final, nous avons exécuté un code séquentiel, exactement le **même code** sans modification (à quelques déclarations de variables près) sur des processeurs séparés de 30 ans sur des systèmes d'exploitations entre 1996 (la Debian Buzz) et 2019 (la Ubuntu 18.10 ou la Debian Buster).
+Au final, nous avons exécuté un code séquentiel, exactement le **même code** sans modification (à quelques déclarations de variables près) sur des processeurs séparés de 30 ans sur des systèmes d'exploitation entre 1996 (la Debian Buzz) et 2019 (la Ubuntu 18.10 ou la Debian Buster).
 Le gain sans avoir modifié le code dépasse (ou approche) le million : 3 millions en BB mais "seulement" 700000 en Mono séparent le Intel 80386SX du Threadripper 1950X.
@@ Ligne 226: / Ligne 224: @@
 {{:developpement:activites:qualification:serialallprocessorsvsfreq.png?500|}}
-Nous constatons que les processeurs au sein d'une même génération sont complètement comparables (Skylake avec les Silver 4114, Gold 5122 et W-2145) en Mono et en BB. Il n'y a pas, en 15 ans (depuis le P4 Northwood) un gain supérieur à 3 en Mono. Le gain en BB est beaucoup plus significatif sur cette période : autour de 15.
+Nous constatons que les processeurs au sein d'une même génération sont parfaitement comparables (Skylake avec les Silver 4114, Gold 5122 et W-2145) en Mono et en BB. Il n'y a pas, en 15 ans (depuis le P4 Northwood) un gain supérieur à 3 en Mono. Le gain en BB est beaucoup plus significatif sur cette période : autour de 15.
 Ainsi, le processeur le plus rapide sur notre problème est 3430x plus rapide en BB et plus de 7763x en Mono qu'un 80386SX à fréquence équivalente.
 === Et la loi de Moore dans tout ça ? ===
@@ Ligne 239: / Ligne 237: @@
   * une période entre 2000 et 2019 de faible croissance
-Entre 1994 et 2000, les deux droites de régression permettent d'estimer un doublement de performances au bout d'une année seulement (1.01 et 1.1 années en BB et en Mono respectivement) : l'âge d'or de la microinformatique !
+Entre 1994 et 2000, les deux droites de régression permettent d'estimer un doublement de performances au bout d'une année seulement (1.01 et 1.1 années en BB et en Mono respectivement) : l'âge d'or de la micro-informatique !
 Mais à partir de 2004, un vrai tassement avec "presque" un pallier de performances comme le montrait le schéma précédent : en fait, seule la performance en BB augmente significativement d'une génération de processeur à la suivante (20x entre le ThreadRipper et le Northwood en BB et seulement 4x en Mono).
@@ Ligne 247: / Ligne 245: @@
 ===== La parallélisation du code =====
-Nous allons explorer plusieurs approches de parallélisation. Tout d'abord, la première va chercher à minimiser au maximum la modification du code source avec OpenMP. Puis, la seconde va explorer les approches de distribution massives de tâches avec OpenCL sur CPU : cette méthode permettra en outre, par son universalité, de pouvoir "aussi" être exploitée sur les GPU. Enfin, cette approche OpenCL sera modifier pour exploiter la méthode classique d'exploitation des GPU Nvidia via CUDA.
+Nous allons explorer plusieurs approches de parallélisation. Tout d'abord, la première va chercher à minimiser au maximum la modification du code source avec OpenMP. Puis, la seconde va explorer les approches de distribution massives de tâches avec OpenCL sur CPU : cette méthode permettra en outre, par son universalité, de pouvoir "aussi" être exploitée sur les GPU. Enfin, cette approche OpenCL sera modifiée pour exploiter la méthode classique d'exploitation des GPU Nvidia via CUDA.
 ==== Parallélisation à la OpenMP ====
@@ Ligne 253: / Ligne 251: @@
 Programmer avec OpenMP se résume souvent de la manière suivante : "casser des boucles". En fait, il s'agit plutôt de respecter le premier principe de la parallélisation : on distribue ce qui peut l'être le plus longtemps possible.
-Dans notre cas, la génération de notre image par //lancer de rayons// ne s'est pas faite pixel par pixel avec un calcul individuel des trajectoires. Nous avons exploité la symétrie du problème : la trajectoire ne dépend "que" de la distance qui sépare l'axe de l'image (pointant le centre du trou noir) avec la direction initiale du rayon partant de notre image. Cette valeur est le "paramètre d'impact". Ensuite, un fois cette trajectoire calculée pour ce paramètre d'impact, tous les pixels sur ce cercle étaient explorés pour savoir si oui ou non ils correspondaient à des trajectoires entrant en collision avec le disque.
+Dans notre cas, la génération de notre image par //lancer de rayons// ne s'est pas faite pixel par pixel avec un calcul individuel des trajectoires. Nous avons exploité la symétrie du problème : la trajectoire ne dépend "que" de la distance qui sépare l'axe de l'image (pointant le centre du trou noir) avec la direction initiale du rayon partant de notre image. Cette valeur est le "paramètre d'impact". Ensuite, une fois cette trajectoire calculée pour ce paramètre d'impact, tous les pixels sur ce cercle étaient explorés pour savoir si oui ou non ils correspondaient à des trajectoires entrant en collision avec le disque.
 La parallélisation naturelle s'est donc développée sur ce paramètre d'impact : pour une image 256x256, il y a 128 paramètres d'impact à explorer, lesquels peuvent être traités indépendamment, donc parallélisés. Cependant, nous pressentons un petit souci : au fur et à mesure que nous nous éloignons du centre de l'image, le cercle à explorer grandit et donc le nombre de pixels à explorer aussi. Traiter donc un paramètre d'impact "faible" prendra beaucoup moins de temps qu'un cercle de paramètre d'impact sur le bord de l'image.
-La modification initiale du programme est plutôt simple : rajouter quelques directives (préfixées par ''#''). Cependant, il faut quand même prendre quelques précautions avec le code : le fait que chaque trajectoire soit traitée indépendamment impose de placer les déclarations de variables exploitées après les directives indiquant la portion à paralléliser.
+La modification initiale du programme est plutôt simple : rajouter quelques directives (préfixées par ''#''). Cependant, il faut quand même prendre quelques précautions avec le code : le fait que chaque trajectoire est traitée indépendamment impose de placer les déclarations de variables exploitées après les directives indiquant la portion à paralléliser.
-Dernier détail : lors de l'exécution d'un programme parallélisé avec OpenMP, le système d'exploitation "détecte" le nombre d'unités de calcul, le nombre de "coeurs" puis parallélise sur ce nombre de coeurs, c'est à dire lance un nombre de tâches simultanées à concurrence du nombre de coeurs disponibles. Il est aussi possible, via la variable d'environnement du système ''OMP_NUM_THREADS'', de "contrôler" le nombre de tâches qui seront exécutées simultanément.
+Dernier détail : lors de l'exécution d'un programme parallélisé avec OpenMP, le système d'exploitation "détecte" le nombre d'unités de calcul (le nombre de "coeurs") puis parallélise sur ce nombre de coeurs, c'est à dire lance un nombre de tâches simultanées à concurrence du nombre de coeurs disponibles. Il est aussi possible, via la variable d'environnement du système ''OMP_NUM_THREADS'', de "contrôler" le nombre de tâches qui seront exécutées simultanément.
-Il est alors intéressant, avec les commandes "système", de contrôler si tous les coeurs sont bien exploités efficacement.
+Il est donc intéressant, avec les commandes "système", de contrôler si tous les coeurs sont bien exploités efficacement.
 Pour ne pas surcharger les résultats, nous séparons les simulations BB et Mono.
@@ Ligne 267: / Ligne 265: @@
 {{ :developpement:activites:qualification:openmp_bb.png?500 |}}
-En bleu apparaissent les performances du code séquentiel en BB, non parallélisé, déjà présentés... Effectivement, le résultat est impressionnant. Les deux premiers processeurs à être équipés multi-coeurs logiques ou physiques sont invisibles. Nous constatons cependant que les machines équipées de 8 coeurs physiques et relativement récents ont des performances comparables (W-2145, Gold 5122x2 et E5-2637v4x2). Les machines équipées de coeurs plus anciens présentent des performances inférieures à la moitié voire au cinquième. Autre logique respectée, les machines disposant de plus de 8 coeurs physiques (16, 20 et même 28 coeurs) respectent la logique : elles forment le podium.
+En bleu apparaissent les performances du code séquentiel en BB, non parallélisé, déjà présentées... Effectivement, le résultat est impressionnant. Les deux premiers processeurs à être équipés de multi-coeurs logiques ou physiques sont invisibles. Nous constatons cependant que les machines équipées de 8 coeurs physiques et relativement récents ont des performances comparables (W-2145, Gold 5122x2 et E5-2637v4x2). Les machines équipées de coeurs plus anciens présentent des performances inférieures à la moitié voire au cinquième. Autre logique respectée, les machines disposant de plus de 8 coeurs physiques (16, 20 et même 28 coeurs) respectent la logique : elles forment le podium.
 Regardons s'il en est de même en Mono.
@@ Ligne 273: / Ligne 271: @@
 {{ :developpement:activites:qualification:openmp_mono.png?500 |}}
-La même distribution apparaît, avec des ratio de performances encore plus marqués par rapport au code séquentiel.
+La même distribution apparaît, avec des ratios de performances encore plus marqués par rapport au code séquentiel.
 Ainsi, le parallélisme en OpenMP permet de porter notre gain en performance de 3 millions à 68 millions en B  et de 700000 à plus de 14 millions : un gros facteur 20 ! De quoi rattraper la loi de Moore largement maltraitée par le code purement séquentiel.
-En fait, pour être honnête, ces résultats n'ont pu être obtenus sans modification du "système". En effet, lors de l'exécution, il est apparu que tous les coeurs n'étaient pas sollicités de manière équivalentes et donc que l'exécution n'était pas optimale. Pour obtenir de telles performances, nous avons littéralement "bourrer" les coeurs avec beaucoup plus de tâches simultanées qu'ils n'en existaient de physiques : jusqu'à 128x le nombre de coeurs physiques.
+En fait, pour être honnête, ces résultats n'ont pu être obtenus sans modification du "système". En effet, lors de l'exécution, il est apparu que tous les coeurs n'étaient pas sollicités de manière équivalente et donc que l'exécution n'était pas optimale. Pour obtenir de telles performances, nous avons littéralement "bourré" les coeurs avec beaucoup plus de tâches simultanées qu'il n'en existait de physiques : jusqu'à 128x le nombre de coeurs physiques.
 Pour illustrer cela, nous avons exploré le temps de calcul pour différentes tailles d'images et nous avons joué sur le parallélisme imposé au code OpenMP, allant de 1x le nombre de coeurs à 128x le nombre de coeurs, pour les deux méthodes Mono et BB.
-La première représentation parle d'elle même, en appliquant ces paramètres sur la méthode BB.
+La première représentation parle d'elle-même, en appliquant ces paramètres sur la méthode BB.
 {{ :developpement:activites:qualification:tr_omp_bb.png?500 |}}
@@ Ligne 291: / Ligne 289: @@
 {{ :developpement:activites:qualification:tr_omp_mono.png?500 |}}
-Ici, c'est encore pire ! Pour un parallélisme fixé à 16 tâches, nous n'avons qu'un ratio de 2. Par contre, au delà de 64x le nombre de coeurs, la performance en OpenMP dépasse 16x la performance en séquentiel.
+Ici, c'est encore pire ! Pour un parallélisme fixé à 16 tâches, nous n'avons qu'un ratio de 2. Par contre, au-delà de 64x le nombre de coeurs, la performance en OpenMP dépasse 16x la performance en séquentiel.
 Il existe certainement des méthodes plus "astucieuses" pour faire cela directement à l'intérieur du code source. Mais nous voulions illustrer que l'exploitation optimale des performances ne passe pas "que" par le code source. Cela passe par les options de compilation (que nous n'avons pas évoquée) mais surtout par une connaissance du système d'exploitation et de ses mécanismes.
@@ Ligne 297: / Ligne 295: @@
 ==== Parallélisation avec OpenCL ====
-Bien qu'existant depuis plus de 10 ans [[https://fr.wikipedia.org/wiki/OpenCL|OpenCL]] est un mal aimé dans les approches de programmation parallèles. Et pourtant, il dispose d'un atout de choix : son universalité  ; il est en effet possible, en programmant "bien" en OpenCL, d'avoir un même programme fonctionnant sans modification et sans recompilation, sur la presque quasi-totalité des GPU, sur des FPGA Intel ou Xilinx, sur des DSP mais aussi sur les processeurs.
+Bien qu'existant depuis plus de 10 ans [[https://fr.wikipedia.org/wiki/OpenCL|OpenCL]] est un mal-aimé dans les approches de programmation parallèle. Et pourtant, il dispose d'un atout de choix : son universalité  ; il est en effet possible, en programmant "bien" en OpenCL, d'avoir un même programme fonctionnant sans modification et sans recompilation, sur la presque quasi-totalité des GPU, sur des FPGA Intel ou Xilinx, sur des DSP mais aussi sur les processeurs.
-Il existe ainsi 3 implémentations OpenCL pour les processeurs : l'historique et la première développée par AMD, déjà vieillissante mais exploitable sur les anciens processeurs (avant 2009), l'Open Source et l'inefficace dans Portable CL, dont le seul mérite est d'exister et l'implémentation Intel, méconnue et bien suivie (mises à jour régulières).
+Il existe ainsi 3 implémentations OpenCL pour les processeurs : l'historique et la première développée par AMD, déjà vieillissante mais exploitable sur les anciens processeurs (avant 2009), l'Open Source et l'inéfficace Portable CL, dont le seul mérite est d'exister et l'implémentation Intel, méconnue et bien suivie (mises à jour régulières).
-Le portage de notre code en OpenCL va cependant être plus long que sa parallélisation en OpenMP. Il existe des approches comparables avec OpenACC, fonctionnelles, mais elles exigent soit l'exploitation du compilateur PGI (cher pour un non académique), soit l'exploitation d'un compilateur effroyablement récent.
+Le portage de notre code en OpenCL va cependant être plus long que sa parallélisation en OpenMP. Il existe également des approches comparables avec OpenACC, fonctionnelles, mais elles exigent soit l'exploitation du compilateur PGI (cher pour un non académique), soit l'exploitation d'un compilateur "trop" récent.
-Précisons tout de suite que programmer en OpenCL //from scratch// dans du C pur est effroyable : par exemple, il faut intégrer DANS le code source le code OpenCL et préciser quelques lignes plus bas le nombre de lignes de code dans un noyau OpenCL. En effet, tout l'intérêt de OpenCL va être d'isoler les parties du code à paralléliser dans des "noyaux" lesquels seront appelés en parallèle, massivement, très massivement. Pour GPU récent, il est possible de montrer que l'optimum s'obtient pour au moins 32 fois le nombre d'unités de calcul. Si nous avons un processeur avec 8 coeurs, ce sera 256 fois. Si nous avons un GPU à 4352 coeurs, ce sera presque 140000. Tout le travail donc être d'extraire du code source les parties qui feront partie de chacun de ces "noyaux".
+Précisons tout de suite que programmer en OpenCL //from scratch// dans du C pur est effroyable : par exemple, il faut intégrer DANS le code source le code OpenCL et préciser quelques lignes plus bas le nombre de lignes de code dans un noyau OpenCL. En effet, tout l'intérêt de OpenCL va être d'isoler les parties du code à paralléliser dans des "noyaux", lesquels seront appelés en parallèle, massivement, très massivement. Pour GPU récent, il est possible de montrer que l'optimum s'obtient pour au moins 32 fois le nombre d'unités de calcul. Si nous avons un processeur avec 8 coeurs, ce sera 256 fois. Si nous avons un GPU à 4352 coeurs, ce sera presque 140000. Tout le travail donc être d'extraire du code source les parties qui feront partie de chacun de ces "noyaux".
-Plutôt que de coder en C ou en C++, beaucoup trop difficile à maintenir, nous choisissons l'option Python. En effet, [[https://documen.tician.de/pyopencl/|PyOpenCL]] existe depuis plus de 10 ans et permet de simplifier dramatiquement l'exploitation de OpenCL. De plus, il existe également le pendant CUDA avec [[https://documen.tician.de/pycuda/|PyCUDA]]. Profitons en pour remercier Andreas Kloeckner pour ces superbes outils !
+Plutôt que de coder en C ou en C++, beaucoup trop difficiles à maintenir, nous choisissons l'option Python. En effet, [[https://documen.tician.de/pyopencl/|PyOpenCL]] existe depuis plus de 10 ans et permet de simplifier dramatiquement l'exploitation de OpenCL. De plus, il existe également le pendant CUDA avec [[https://documen.tician.de/pycuda/|PyCUDA]]. Profitons-en pour remercier Andreas Kloeckner pour ces superbes outils !
-Sachant que, pour exploiter au mieux nos machines, il faut paralléliser au maximum, regardons comme distribuer nos tâches en OpenCL. La première, c'est le code sans la moindre parallélisation, ce sera **Original**. La seconde, ce sera d'appliquer la même méthode qu'en OpenMP : ce sera **EachCircle** où chaque tâche aura à traiter un paramètre d'impact, donc parallélisme équivalent à la moitié de la taille de l'image. Le troisième, ce sera d'exploiter du //raytracing// brutal, où, pour chaque pixel, la trajectoire est recherchée : ce sera la méthode **EachPixel**. Les deux dernières méthodes seront des méthodes hybrides : la première calculera toutes les trajectoires en parallèle, puis la seconde estimera en parallèle pour tous les secteurs de chaque cercle associé et chaque paramètre d'impact la valeur de chaque pixel ; ce sera **TrajectoCircle**. La dernière exploitera les trajectoires calculées dans une première phase dans une seconde en associant directement un trajectoire à un pixel donné.
+Sachant que, pour exploiter au mieux nos machines, il faut paralléliser au maximum, regardons comment distribuer nos tâches en OpenCL. La première, c'est le code sans la moindre parallélisation, ce sera **Original**. La deuxième, ce sera d'appliquer la même méthode qu'en OpenMP : ce sera **EachCircle** où chaque tâche aura à traiter un paramètre d'impact, donc  un parallélisme équivalent à la moitié de la taille de l'image. La troisième, ce sera d'exploiter du //raytracing// brutal, où, pour chaque pixel, la trajectoire est recherchée : ce sera la méthode **EachPixel**. Les deux dernières méthodes seront des méthodes hybrides : la première calculera toutes les trajectoires en parallèle, puis la seconde estimera en parallèle, pour tous les secteurs de chaque cercle, et chaque paramètre d'impact la valeur de chaque pixel ; ce sera **TrajectoCircle**. La dernière exploitera les trajectoires calculées dans une première phase, puis en une seconde, en associant directement une trajectoire à un pixel donné.
-Ainsi, nous pouvons résumer, pour une image de taille NxN, le nombre de tâches exécutables en parallèles pour chaque méthode, le "régime de parallélisme" en somme :
+Ainsi, nous pouvons résumer, pour une image de taille NxN, le nombre de tâches exécutables en parallèle pour chaque méthode, le "régime de parallélisme" en somme :
   * **Original** : 1
   * **EachCircle** : N/2
@@ Ligne 318: / Ligne 316: @@
 Chacune des méthodes (à l'exception de **Original**) sera donc étudiée et comparée aux autres pour des images allant de 64x64 à 16384x16384. Dix lancements successifs pour chaque méthode permettront ensuite de disposer d'une statistique pertinente.
-Petit détail pour la métrologie : dans le cas du programme originel, le temps de calcul était estimé à partir de plusieurs //timers// : l'un en secondes indexé sur l'horloge temps réel (''time(NULL)''), l'un en microsecondes sur la même horloge (''getttimeofday()''), le dernier exploitant ''cputime()''. Avec OpenCL, le principe est d'exploiter un "périphérique", même si ce dernier est en fait le processeur. Les pesaces mémoire exploités doivent être échangés entre l'hôte et le périphérique, et cela prend du temps ! Ainsi, il faut séparer deux durées : celle du calcul proprement dit (ce sera **Compute**), et celle du "temps écoulé", ajoutant le temps de calcul au temps de transfert de la mémoire du périphérique à l'hôte (ce sera le **Elapsed**).
+Petit détail pour la métrologie : dans le cas du programme originel, le temps de calcul était estimé à partir de plusieurs //timers// : l'un en secondes indexé sur l'horloge temps réel (''time(NULL)''), l'un en microsecondes sur la même horloge (''getttimeofday()''), le dernier exploitant ''cputime()''. Avec OpenCL, le principe est d'exploiter un "périphérique", même si ce dernier est en fait le processeur. Les espaces "mémoire" exploités doivent être échangés entre l'hôte et le périphérique, et cela prend du temps ! Ainsi, il faut séparer deux durées : celle du calcul proprement dit (ce sera **Compute**), et celle du "temps écoulé", ajoutant le temps de calcul au temps de transfert de la mémoire du périphérique à l'hôte (ce sera le **Elapsed**).
 === Premier lancement en OpenCL ===
-Illustrons cela sur le processeur ''TOP 1'' de notre exécution séquentielle, sans parallélisation : sur le ThreadRipper 1950X.  Et comparons donc les temps de calcul ou écoulé pour les 7 méthodes pour une grosse image 16384x16384 : **Serial** sans parallélisation, **OpenMP** et les 5 méthodes exploitant OpenCL.
+Illustrons cela sur le processeur ''TOP 1'' de notre exécution séquentielle, sans parallélisation : sur le ThreadRipper 1950X.  Et comparons donc les temps de calcul ou écoulés, en employant les 7 méthodes, pour une grosse image 16384x16384 : **Serial** sans parallélisation, **OpenMP** et les 5 méthodes exploitant OpenCL.
 {{ :developpement:activites:qualification:tr1950x_7methods_elapsed.png?500 |}}
@@ Ligne 342: / Ligne 340: @@
 Inutile de préciser que, pour les processeurs antérieurs à 2008, leurs performances sont imperceptibles.
-En Mono, nous retrouvons le facteur 2 de performances entre OpenCL et OpenMP pour le ThreadRipper. Mais pour les processeurs Intel récents, c'est encore meilleur ! Les derniers Skylake d'Intel (W-2145 et Gold 5122) présentent des performances entre 4 et 5 fois meilleures qu'en OpenMP ! Etrangement, le Silver 4114, pourtant Skylake, ne gagne "que" 70% mais le numéro 1, le système disposant de 28 coeurs à base de E5-2680v4, double ses performances. Tous les processeurs Intel gagnent un OpenCL, même l'antique système à base de E5440 ne fonctionnant qu'avec l'implémentation OpenCL d'AMD.
+En Mono, nous retrouvons un facteur 2 de performances entre OpenCL et OpenMP pour le ThreadRipper. Mais pour les processeurs Intel récents, c'est bien meilleur ! Les derniers Skylake d'Intel (W-2145 et Gold 5122) présentent des performances entre 4 et 5 fois meilleures qu'en OpenMP ! Etrangement, le Silver 4114, pourtant Skylake, ne gagne "que" 70% mais le numéro 1, le système disposant de 28 coeurs à base de E5-2680v4, double ses performances. Tous les processeurs Intel gagnent un OpenCL, même l'antique système à base de E5440 ne fonctionnant qu'avec l'implémentation OpenCL d'AMD.
-En BB, le constat est plus nuancé. Pour le Threadripper, les performances OpenCL et OpenMP sont équivalentes et rétrograde de la seconde à la cinquième place. Les processeurs Intel s'en sortent plutôt mieux au point que le dernier Skylake W-2145 à 8 coeurs détrône le système à E5-2680v4 à 28 coeurs ! Pour les processeurs anciens X5550 et E5440, l'OpenMP fait bien mieux.
+En BB, le constat est plus nuancé. Pour le Threadripper, les performances OpenCL et OpenMP sont équivalentes et rétrogradent de la seconde à la cinquième place. Les processeurs Intel s'en sortent plutôt mieux au point que le dernier Skylake W-2145 à 8 coeurs détrône le système à base de E5-2680v4 avec 28 coeurs ! Pour les processeurs anciens X5550 et E5440, l'OpenMP fait bien mieux...
 Regardons ces performances maintenant en intégrant les temps de transferts
@@ Ligne 354: / Ligne 352: @@
 {{ :developpement:activites:qualification:cpu_elapsed_bb.png?500 |}}
-En BB, nous réalisons quasiment les mêmes constats : Threadripper de second à cinquième, meilleur en OpenMP que OpenCL ; les Skylake bien meilleurs d'un facteur presque 3 en OpenCL ; le Silver et les deux Broadwell avec quelques % de mieux. Le système E5-2680v4
+En BB, nous réalisons quasiment les mêmes constats : Threadripper de deuxième à cinquième, meilleur en OpenMP que OpenCL ; les Skylake bien meilleurs d'un facteur presque 3 en OpenCL ; le Silver et les deux Broadwell avec quelques % de mieux.
 Quelles conclusions tirer de ces comparaisons ?
-Tout d'abord qu'Intel a très bien travaillé son OpenCL pour ses propres processeurs... Sans les Skylake, essentiellement le Gold et le W-2145, les gains seraient substantiels pour la partie **Compute** mais fondent dès que les temps de transfert sont intégrés. Puis que le Threadripper "résiste", mais pas pour tous les usages : un test trop sommaire peut donc rapidement induire en erreur quant à la performance d'un processeur. Ensuite que les processeurs anciens, dès lors que l'implémentation Intel est exploitable (donc après la génération Nehalem), l'OpenCL reste une bonne option. Enfin que, avec OpenCL, notre ratio de performances de calcul frise les 100 millions en BB et 35 millions en Mono. En intégrant les transferts, ces ratios tombent à 81 millions en BB et 17 millions en Mono : 8 ordres de grandeurs séparent des processeurs séparés de 30 ans.
+Tout d'abord qu'Intel a très bien travaillé son OpenCL pour ses propres processeurs... Sans les Skylake, essentiellement le Gold et le W-2145, les gains seraient substantiels pour la partie **Compute** mais fondent dès que les temps de transfert sont intégrés.
+Puis que le Threadripper "résiste", mais pas pour tous les usages : un test trop sommaire peut donc rapidement induire en erreur quant à la performance d'un processeur.
+Ensuite que les processeurs anciens, dès lors que l'implémentation Intel est exploitable (donc après la génération Nehalem), l'OpenCL reste une bonne option.
+Enfin que, avec OpenCL, notre ratio de performances de calcul frise les 100 millions en BB et 35 millions en Mono. En intégrant les transferts, ces ratios tombent à 81 millions en BB et 17 millions en Mono : 8 ordres de grandeurs distinguent des processeurs séparés de 30 ans.
 === Intégration des accélérateurs en OpenCL ===
@@ Ligne 370: / Ligne 374: @@
 {{ :developpement:activites:qualification:c1060_harpertown.png?500 |}}
-En Mono, la Tesla C1060 est 3.5x plus rapide que le serveur bi-socket en OpenCL, l'OpenMP étant très légèrement inférieur en performances. En BB, c'est une toute autre affaire : l'OpenCL de la Tesla est à peine supérieure à la version OpenMP du serveur bi-socket. La version OpenCL d'AMD, exploitée ici, est dramatiquement inefficace face à l'OpenMP. Ainsi, déjà, nous constatons que, poiur notre première exploitation de GPU, son comportement n'est pas du tout éauivalent en fonction de la "charge" que nous portons à chaque processus élémentaire.
+En Mono, la Tesla C1060 est 3.5x plus rapide que le serveur bi-sockets en OpenCL, l'OpenMP étant très légèrement inférieur en performances. En BB, c'est une toute autre affaire : l'OpenCL de la Tesla est à peine supérieure à la version OpenMP du serveur bi-socket. La version OpenCL d'AMD, exploitée ici, est dramatiquement inefficace face à l'OpenMP. Ainsi, déjà, nous constatons que, pour notre première exploitation de GPU, son comportement n'est pas du tout équivalent en fonction de la "charge" que nous portons à chaque processus élémentaire.
 == AMD : les trois dernières générations ==
@@ Ligne 378: / Ligne 382: @@
 {{ :developpement:activites:qualification:amdvsbestcpu.png?500 |}}
-Nous constatons que ce n'est "que" pour la dernière génération de GPU que AMD dépasse sensiblement nos meilleurs CPU. Avec la Radeon 7, les performances sont entre 2 et 3 fois supérieures en BB et en Mono. Le ratio de performances entre Mono et BB passe de presque 10 à moins de 4, preuve que AMD, dans ses dernières générations adopte un comportement à la charge comparable à celui d'un processeur.
+Nous constatons que ce n'est "que", pour la dernière génération de GPU, que AMD dépasse sensiblement nos meilleurs CPU. Avec la Radeon 7, les performances sont entre 2 et 3 fois supérieures en BB et en Mono. Le ratio de performances entre Mono et BB passe de presque 10 à moins de 4, preuve que AMD, dans ses dernières générations adopte un comportement "à la charge" comparable à celui d'un processeur.
 == Nvidia : 5 générations de "Gaming" ==
-A l'origine, c'est bien le détournement des cartes de "gaming" grand initateur de l'exploitation des GPU en informatique scientifique. Il est donc naturel de comparer nos deux meilleurs processeurs à 5 générations successives de cartes Nvidia en OpenCL : la GTX560Ti (circuit Fermi), la GTX780Ti (circuit Kepler), la GTX980Ti (circuit Maxwell), la GTX 1080 Ti (circuit Pascal) et la RTX Titan (circuit Turing).
+A l'origine, c'est bien le détournement des cartes de "gaming" grand initiateur de l'exploitation des GPU en informatique scientifique. Il est donc naturel de comparer nos deux meilleurs processeurs à 5 générations successives de cartes Nvidia en OpenCL : la GTX560Ti (circuit Fermi), la GTX780Ti (circuit Kepler), la GTX980Ti (circuit Maxwell), la GTX 1080 Ti (circuit Pascal) et la RTX Titan (circuit Turing).
 {{ :developpement:activites:qualification:nvidiagamervsbestcpu.png?500 |}}
-En Mono, nous constatons que, dès la GTX780Ti (de fin 2013), un GPU de Gamer dépasse n'importe lequel des processeur. La progression en Mono des performances et (presque linéaire) entre Fermi et Turing, à l'exception de la génération Maxwell, un peu en deça des performances attendues. En BB, c'est bien différent : il faut attendre la toute dernière génération pour obtenir une performance comparable. Il existe toujours également un ratio spectaculaire de performances pour les GPU de Gamer entre Mono et BB : largement supérieur à 10.
+En Mono, nous constatons que, dès la GTX780Ti (de fin 2013), un GPU de Gamer dépasse n'importe lequel des processeurs. La progression en Mono des performances et (presque linéaire) entre Fermi et Turing, à l'exception de la génération Maxwell, un peu en deça des performances attendues. En BB, c'est bien différent : il faut attendre la toute dernière génération pour obtenir une performance comparable. Il existe toujours également un ratio spectaculaire de performances pour les GPU de Gamer entre Mono et BB : largement supérieur à 10.
 == Et les accélérateurs non GPU dans tout ça ? ==
@@ Ligne 418: / Ligne 422: @@
 Le portable inital du code en CUDA a été plutôt rapide : il n'a pris qu'une soirée. Il a suffit de remplacer tous les "workitem" par "BlockIdx" et préfixer les fonctions primaires par  "__device__".
-Mais la première exécution a été plutôt très décevante : en effet, sur la Tesla P100, j'avais un temps de référence de 1 seconde d'exécution en OpenCL sur une simulation BB. En CUDA, ce temps est devenu 24 secondes (soit 24x plus lent !). En Mono, le temps était un peu supérieur.
+Mais la première exécution a été plutôt très décevante : en effet, sur la Tesla P100, nous avions un temps de référence de 1 seconde d'exécution en OpenCL sur une simulation BB. En CUDA, ce temps est devenu 24 secondes (soit 24x plus lent !). En Mono, le temps était un peu supérieur.
 == Exploiter les deux étages de parallélisme : une nécessité mais comment ? ==
@@ Ligne 424: / Ligne 428: @@
 En effet, pour exploiter efficacement CUDA, il faut IMPERATIVEMENT solliciter les "Threads". En effet, la programmation en CUDA (et en OpenCL) désarçonne le néophyte en présentant deux "étages" de parallélisme. Si l'existence de ces deux étages est pleinement compréhensible au regard d'une multiplication de matrices (chaque élément de la matrice résultante est indépendant de tous les autres mais le calcul de chaque élément lui aussi peut être parallélisé), elle impose une gestion plus compliquée de la distribution des tâches, essentiellement parce que, si le premier étage permet de lancer plus d'un milliard de tâches concurrentes, le second ne permet d'en lancer que 1024.
-Dans notre cas, il ne suffit plus de remplacer un //workitem// ''get_global_id(0)'' définissant l'indice d'une tâche élémentaure par un ''BlockIdx.x'' mais par une combinaison finement choisie de ''BlockIdx'' et ''ThreadIdx''. Ce qui donne, dans notre cas, pour une position ''xi'', nous avons en OpenCL :
+Dans notre cas, il ne suffit plus de remplacer un //workitem// ''get_global_id(0)'' définissant l'indice d'une tâche élémentaire par un ''BlockIdx.x'' mais par une combinaison finement choisie de ''BlockIdx'' et ''ThreadIdx''. Dans notre cas, pour une position ''xi'', nous avons en OpenCL :
 <code>
 uint xi=(uint)get_global_id(0);
@@ Ligne 439: / Ligne 443: @@
 == Exploitation de CUDA : tirer le maximum, mais des plus récents ==
-Commençons par la simulation en BB : nous constatons tout d'abord que certaines architectures ne fonctionnent pas en CUDA avec notre code. En effet, les machines les plus anciennes (GTX 560 Ti, Tesla C1060, Teska M2090) plantent à l'exécution de CUDA. En effet, le CUDA version 10.1 exploité ne supporte pas les anciens pilotes 340, les derniers disponibles dans la distributions permettant le support de ces GPU ou GPGOU vieillissant.
+Commençons par la simulation en BB : nous constatons tout d'abord que certaines architectures ne fonctionnent pas en CUDA avec notre code. En effet, les machines les plus anciennes (GTX 560 Ti, Tesla C1060, Tesla M2090) plantent à l'exécution de CUDA. En effet, le CUDA version 10.1 exploité ne supporte pas les anciens pilotes 340, les derniers disponibles dans la distribution Debian permettant le support de ces GPU ou GPGPU vieillissants.
 {{ :developpement:activites:qualification:cudavsall_compute_bb.png?500 |}}
-Nous remarquons que seules les cartes Tesla récentes présentent des accélérations notables au passage en CUDA. La Tesla V100 est plus de deux fois plus rapide, la P100 2/3 plus rapide. La dynamique effondre donc tous les autres accélérateurs et plus encore processeurs. Pour les GPU de //gamer//, CUDA fait à peine mieux que OpenCL. Il se passe donc des "choses bizarres" lors de la phase de compilation des noyaux CUDA avec ''NVCC''. Dès la Tesla K40m, nous sont plus performants que n'importe quel autre accélérateur. Quant à la Tesla V100, elle règle sans partage, reléguant la P100 à un facteur 2.
+Nous remarquons que seules les cartes Tesla récentes présentent des accélérations notables au passage en CUDA. La Tesla V100 est  deux fois plus rapide, la P100 2/3 plus rapide. La dynamique effondre donc tous les autres accélérateurs et plus encore processeurs. Pour les GPU de //gamer//, CUDA fait à peine mieux que OpenCL. Il se passe donc des "choses bizarres" lors de la phase de compilation des noyaux CUDA avec ''NVCC''. Dès la Tesla K40m, les GPGPU sont plus performants que n'importe quel autre accélérateur de gamer. Quant à la Tesla V100, elle règne sans partage, reléguant la P100 à un facteur 2.
 {{ :developpement:activites:qualification:cudavsall_compute_mono.png?500 |}}
-Pour une simulation Mono, les gains sont beaucoup plus significatifs. Pour les GTX ou RTX, les performances quintuplent. Pour les Tesla, nous dépassons le décuplement. La Tesla V100 est moins dominatrice que précédemment face aux RTX ou GTX, mais elle place son facteur 2. Face aux processeurs ou même tous les autres GPU, il y a la Tesla V100 et les autres.
+Pour une simulation Mono, les gains sont beaucoup plus significatifs. Pour les GTX ou RTX, les performances quintuplent. Pour les Tesla, nous dépassons le décuplement. La Tesla V100 est moins dominatrice que précédemment face aux RTX ou GTX, mais elle impose un facteur 2 aux précédentes. Face aux processeurs ou même à tous les autres GPU, il y a la Tesla V100 et les autres.
-Ainsi, l'exploitation CUDA apporte de très substanciels gain en performance. Compte-tenu du temps à passer au passage OpenCL/CUDA, il serait dommage de s'en passer.
+Ainsi, l'exploitation CUDA apporte de très substantiels gains en performance. Compte-tenu du temps passé au portage de OpenCL à CUDA, il serait dommage de s'en priver...
 === Calculer, c'est bien ! Exploiter c'est mieux... ===
-Il ne sert à rien de calculer s'il n'est pas possible d'en exploiter les résultats. Qu'il s'agisse de OpenCL ou CUDA, nous avons vu que tout est périphérique, processeur ou GPU. Il convient donc d'intégrer dans notre simulation le transfert du périphérique vers l'hôte des résultats. Nous allons découvrir que leur importance est loin d'être marginale !
+Il ne sert à rien de calculer s'il n'est pas possible d'en exploiter les résultats. Qu'il s'agisse de OpenCL ou CUDA, nous avons vu que tout est périphérique, processeur ou GPU. Il convient donc d'intégrer dans notre simulation le transfert des résultats du périphérique vers l'hôte. Nous allons découvrir que leur importance est loin d'être marginale !
-L'intégration de ce temps de transfert modifie très sensiblement les classements. Alors que les Nvidia, GPU ou GPGPU, régnaient sans partage avec juste une Radeon 7 plus efficace que les RTX ou GTX en mode BB, la dernière de chez AMD devient comparable à la Tesla P100. Pour une carte 10x moins chère qu'une Tesla V100, elle n'est pas 10x inférieure.
+L'intégration de ce temps de transfert modifie très sensiblement les classements. Alors que les Nvidia, GPU ou GPGPU, régnaient sans partage avec juste une Radeon 7 plus efficace que les RTX ou GTX en mode BB, la dernière de chez AMD devient comparable à la Tesla P100. Pour une carte 10x moins chère qu'une Tesla V100, elle n'est pas 10x inférieure en performances.
 {{ :developpement:activites:qualification:cudavsall_elapsed_bb.png?500 |}}
-En mode Mono, la domination écrasante de la Tesla V100 se réduit à un timide facteur 2 face à la meilleure des AMD. Il est véritablement de voir AMD s'en sortir beaucoup beaucoup mieux que CUDA dans ces transferts, mais peut-être est-ce dû au fait que les 3 cartes AMD présentées ici partagent le même ThreadRipper 1950X comme socle et donc bénéficient de meilleures communications entre hôte et périphériques.
+En mode Mono, la domination écrasante de la Tesla V100 se réduit à un timide facteur 2 face à la meilleure des AMD. Il est véritablement de voir AMD s'en sortir bien mieux que CUDA dans ces transferts, mais peut-être est-ce dû au fait que les 3 cartes AMD présentées ici partagent le même ThreadRipper 1950X comme socle et donc bénéficient de meilleures communications entre hôte et périphériques.
 {{ :developpement:activites:qualification:cudavsall_elapsed_mono.png?500 |}}
@@ Ligne 465: / Ligne 469: @@
 Au final, de 12x face au meilleur du GPU AMD, de 30x face aux meilleurs CPU, la performance du "monstre" V100 se réduit à un facteur entre 4x et 6x pour les processeurs et entre 2x et 3x pour les GPU. Dans ce cas, clairement, le temps de calcul devient marginal au temps de transfert des résultats du GPU à l'hôte.
-Ainsi, nous pouvons, au travers de ces diverses simulations, identifier une troisième frontière. La première est la limite calculatoire (Compute-bound), la seconde est la limite d'échange mémoire (Memory-bound), la troisième est la limite d'échange périphérique (IO-bound).
+Ainsi, nous pouvons, au travers de ces diverses simulations, identifier une troisième frontière. La première est la limite calculatoire (Compute-bound), la deuxième est la limite d'échange mémoire (Memory-bound), la troisième est la limite d'échange périphérique (IO-bound).
 ==== Comparaison finale sur 30 ans de matériel informatique ====
@@ Ligne 471: / Ligne 475: @@
 Nous avions présenté, au début de notre analyse, la distribution des processeurs en fonction de leur année de sortie et de deux paramètres définissant chacun d'eux : leur nombre de transistors et leur fréquence. Cela nous avait permis de conclure que, finalement, nous disposions d'une bonne répartition permettant de revenir à la loi de Moore "statique". Plus loin, nous nous étions rendu compte que, malgré la progression en nombre de transistors, nos performances n'augmentaient pas en conséquence et qu'il fallait modifier le programme pour mieux en exploiter les transistors, répartis sur plusieurs coeurs : la parallélisation du code, d'abord en OpenMP puis en OpenCL, a ainsi permis de largement mieux exploiter ces processeurs récents.
-Temps est venu de comparer tous nos 30 processeurs et accélérateurs, toutes natures confondues, et de les replacer sur un unique graphique, mais en fonction de leur année de commercialisation.
+Le temps est venu de comparer tous nos 30 processeurs et accélérateurs, toutes natures confondues, et de les replacer sur un unique graphique, mais en fonction de leur année de commercialisation.
 {{ :developpement:activites:qualification:performancesyear_bb.png?600 |}}
@@ Ligne 481: / Ligne 485: @@
 Pour les simulations Mono, la différence entre GPU et CPU est beaucoup plus marquée : les triangles verts (CUDA) se détachent beaucoup plus largement. Elément intéressant, ils assurent la continuité de progression en performance assurée, d'abord par les processeurs monocoeur (de 1989 à 2004), puis les multicoeurs (jusqu'à 2008) avant que les GPGPU (en OpenCL) ne proposent des performances supérieures.
 Ainsi, l'exploitation du GPU a eu, dans l'histoire de l'informatique, le même caractère "disruptif" que l'intégration systématique de l'unité en virgule flottante dans les processeurs dès 1993. Grâce à l'exploitation de ces derniers, la loi de Moore "dynamique" (celle sur une performance mesurée) reste encore valable.

developpement/activites/qualification/30ans1code.1576578448.txt.gz · Dernière modification: 2019/12/17 11:27 par equemene

Rechercher

Translations

Navigation

Piste:

Boîte à outils