Différences

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

Lien vers cette vue comparative

Les deux révisions précédentes Révision précédente
Prochaine révision
Révision précédente
Prochaine révision Les deux révisions suivantes
developpement:activites:qualification:30ans1code [2019/11/26 18:03]
equemene [Parallélisation à la OpenMP]
developpement:activites:qualification:30ans1code [2019/12/17 11:27]
equemene
Ligne 14: Ligne 14:
  
 (détournement de Tolkien...) (détournement de Tolkien...)
 +
 ===== Résumé ===== ===== Résumé =====
  
 Partout, on parle d'​augmentation de puissance de calcul des ordinateurs mais comment et en quoi ces capacités ont-elles évolué ces 30 dernières années ? Quoi de plus illustratif que de chercher à "​retrouver"​ par la simulation une image comparable à celle qui a inondé nos médias en ce printemps 2019 : celle d'un trou noir et son écharpe de plasma. Partant du premier article représentant un trou noir, un programme de 25 ans d'âge sera modifié et exécuté sur différents "​moteurs de traitement de l'​information",​ les processeurs,​ classiques, graphiques et éphémères : ceux qui existaient il y a 30 ans, ceux qui émergeaient il y a 10 ans et ceux d'​aujourd'​hui. ​ Partout, on parle d'​augmentation de puissance de calcul des ordinateurs mais comment et en quoi ces capacités ont-elles évolué ces 30 dernières années ? Quoi de plus illustratif que de chercher à "​retrouver"​ par la simulation une image comparable à celle qui a inondé nos médias en ce printemps 2019 : celle d'un trou noir et son écharpe de plasma. Partant du premier article représentant un trou noir, un programme de 25 ans d'âge sera modifié et exécuté sur différents "​moteurs de traitement de l'​information",​ les processeurs,​ classiques, graphiques et éphémères : ceux qui existaient il y a 30 ans, ceux qui émergeaient il y a 10 ans et ceux d'​aujourd'​hui. ​
  
-===== Flash Back de 40 ans : Entre l'​image "​mesurée"​ de 2019 et l'​image simulée de 1979 =====+===== Flash Back de 40 ans : entre l'​image "​mesurée"​ de 2019 et l'​image simulée de 1979 =====
  
 Impossible, en ce printemps 2019, d'​échapper à cette "​image"​ du trou noir de Messier 87. Belle "​image",​ plus dans son sens étymologique que son sens "​photographique"​. Ce qui est dommage, c'est que le traitement médiatique de l'​évènement se soit focalisé sur une brillante programmeuse ayant participé à la réduction de données... Dommage parce que c'est la conjonction du dispositif de synthèse d'​ouverture (qui a bouleversé l'​astronomie ces dernières décennies) et la progression exponentielle de la puissance de calcul scientifique qui ont permis, ensemble, de parvenir à cette prouesse. Impossible, en ce printemps 2019, d'​échapper à cette "​image"​ du trou noir de Messier 87. Belle "​image",​ plus dans son sens étymologique que son sens "​photographique"​. Ce qui est dommage, c'est que le traitement médiatique de l'​évènement se soit focalisé sur une brillante programmeuse ayant participé à la réduction de données... Dommage parce que c'est la conjonction du dispositif de synthèse d'​ouverture (qui a bouleversé l'​astronomie ces dernières décennies) et la progression exponentielle de la puissance de calcul scientifique qui ont permis, ensemble, de parvenir à cette prouesse.
Ligne 24: Ligne 25:
 {{:​developpement:​activites:​qualification:​blackhole20190410.jpg?​400|}} {{:​developpement:​activites:​qualification:​blackhole20190410.jpg?​400|}}
  
-Aussitôt cette "​image"​ publiée, montrant sans ambiguïté l'​horizon du cadavre obscur et son possible disque d'​accrétion en rotation, beaucoup se sont rappelés ​une autre image, la première image (simulée cette fois) d'un disque d'​accrétion autour d'un trou noir. Pour cela, il fallait revenir 40 ans en arrière avec un article publié en 1979 dans Astronomy & Astrophysics par un jeune astrophysicien,​ Jean-Pierre Luminet. Sur la dernière page de son article apparaissait la première image, image issue cette fois d'une expérience numérique. ​+Aussitôt cette "​image"​ publiée, montrant sans ambiguïté l'​horizon du "cadavre obscur" ​et son possible disque d'​accrétion en rotation, beaucoup se sont rappelé ​une autre image, la première image (simulée cette fois) d'un disque d'​accrétion autour d'un trou noir. Pour cela, il fallait revenir 40 ans en arrière avec un article publié en 1979 dans Astronomy & Astrophysics par un jeune astrophysicien,​ Jean-Pierre Luminet. Sur la dernière page de son article apparaissait la première image, image issue cette fois d'une expérience numérique. ​
  
 {{:​developpement:​activites:​qualification:​1979aa_75_228l.png?​400|}} {{:​developpement:​activites:​qualification:​1979aa_75_228l.png?​400|}}
Ligne 40: Ligne 41:
 {{:​developpement:​activites:​qualification:​z_bb.jpg?​300| }} {{:​developpement:​activites:​qualification:​z_bb.jpg?​300| }}
  
-Il aura fallu 22 ans (un cycle solaire ;-) ) pour que ce programme en C serve de matrice à un autre programme permettant d'​exploiter les deux révolutions technologiques majeures de ce 21e siècle en informatique : la multiplication des coeurs et l'​exploitation des GPU comme accélérateur.+Il aura fallu 22 ans (un cycle solaire ;-) ) pour que ce programme en C serve de matrice à un autre programme permettant ​ainsi d'​exploiter les deux révolutions technologiques majeures de ce 21e siècle en informatique : la multiplication des coeurs et l'​exploitation des GPU comme accélérateur.
  
 === Simuler l'​image d'un trou noir, mais comment ? === === Simuler l'​image d'un trou noir, mais comment ? ===
  
-La méthode numérique pour construire cette image est aussi vieille que la génération d'​images par ordinateur : le //​[[https://​fr.wikipedia.org/​wiki/​Ray_tracing|lancer de rayons]]// (ou //​raytracing//​). Le principe consiste à "​remonter le temps" : partir de chaque élément de l'​image (ou de chaque pixel) et suivre son parcours jusqu'​à sa "​source"​. Chaque rayon peut intercepter une surface ou pas. Si non, le pixel est noir, si oui, le pixel a un "​couleur"​ à définir par rapport à son point de contact. Si l'​objet est bleu, le pixel sera bleu mais avec une intensité déterminée à partir ​de des lois physiques. La surface peut aussi être transparente et le rayon peut poursuivre sa route jusqu'​à intercepter une autre surface.+La méthode numérique pour construire cette image est aussi vieille que la génération d'​images par ordinateur : le //​[[https://​fr.wikipedia.org/​wiki/​Ray_tracing|lancer de rayons]]// (ou //​raytracing//​). Le principe consiste à "​remonter le temps" : partir de chaque élément de l'​image (ou de chaque pixel) et suivre son parcours jusqu'​à sa "​source"​. Chaque rayon peut intercepter une surface ou pas. Si non, le pixel est noir, si oui, le pixel a un "​couleur"​ à définir par rapport à son point de contact. Si l'​objet est bleu, le pixel sera bleu mais avec une intensité déterminée à partir des lois physiques. La surface peut aussi être transparente et le rayon peut poursuivre sa route jusqu'​à intercepter une autre surface.
  
 Dans notre cas, les rayons ne sont pas des lignes mais suivent des "​géodésiques de l'​espace-temps",​ lesquelles sont des courbes dans un espace courbe. Il faudra donc estimer, pour chaque pixel, quelle est la trajectoire du rayon. 3 solutions s'​offrent alors pour son origine : Dans notre cas, les rayons ne sont pas des lignes mais suivent des "​géodésiques de l'​espace-temps",​ lesquelles sont des courbes dans un espace courbe. Il faudra donc estimer, pour chaque pixel, quelle est la trajectoire du rayon. 3 solutions s'​offrent alors pour son origine :
-  * soit elle sort de la "​zone"​ de nos objets (trou noir et disque), au delà du rayon externe : le pixel est noir +  * soit elle sort de la "​zone"​ de nos objets (trou noir et disque), au-delà du rayon externe : le pixel est noir ; 
-  * soit elle rentre à un distance inférieure au rayon de Schwarzschild et ne peut en sortir : le pixel est noir (un peu le principe du trou noir, c'est que rien n'en sort en deça du fameux rayon de Schwarzschild) +  * soit elle rentre à un distance inférieure au rayon de Schwarzschild et ne peut en sortir : le pixel est noir (un peu le principe du trou noir, c'est que rien n'en sort en deça du fameux rayon de Schwarzschild) ​; 
-  * soit elle intercepte le disque : le pixel est d'une intensité à définir à partir de l'​intersection entre trajectoire et disque+  * soit elle intercepte le disque : le pixel est d'une intensité à définir à partir de l'​intersection entre trajectoire et disque.
  
 C'est comme cela que toutes les images par lancer de rayons sont construites. C'est comme cela que toutes les images par lancer de rayons sont construites.
Ligne 57: Ligne 58:
 ==== Loi de Moore & loi d'Ohm ==== ==== Loi de Moore & loi d'Ohm ====
  
-En électronique,​ les lois gouvernant le plus l'​évolution technologique sont la loi de Moore et la loi d'​Ohm. ​+En électronique,​ les lois gouvernant le mieux l'​évolution technologique sont la loi de Moore et la loi d'​Ohm. ​
  
-La loi "​commune"​ de Moore stipule que la "​puissance"​ de traitement d'un circuit double tous les 18 mois. En fait, Moore n'a jamais proposé cette loi en l'​état mais a proposé ​qu'un circuit intégrait chaque année deux plus de de transistors pour le même budget que l'​année précédente. Il révisa en 1975 son énoncé en allongeant la période de doublement de transistors à deux années.+La loi "​commune"​ de Moore stipule que la "​puissance"​ de traitement d'un circuit double tous les 18 mois. En fait, Moore n'a jamais proposé cette loi en l'​étatmais a suggéré ​qu'un circuit intégrait chaque année deux plus de transistors pour le même budget que l'​année précédente. Il révisa en 1975 son énoncé en allongeant la période de doublement de transistors à deux années.
  
-La loi d'Ohm décrit elle le comportement endémique d'un composant électrique à dissiper une partie de son énergie sous forme de chaleur quand un courant électrique le traverse. Si cette dernière loi est très utile pour nos bouilloires ou nos machines à laver, elle est plutôt gếnante lorsqu'​il s'agit de composants électroniques,​ notamment les circuits électroniques présents dans tous nos équipements quotidiens.+La loi d'Ohm décritellele comportement endémique d'un composant électrique à dissiper une partie de son énergie sous forme de chaleur quand un courant électrique le traverse. Si cette dernière loi est très utile pour nos bouilloires ou nos machines à laver, elle est plutôt gếnante lorsqu'​il s'agit de composants électroniques,​ notamment les circuits électroniques présents dans tous nos équipements quotidiens.
  
 ==== 30 "​systèmes informatiques"​ et un seul programme ==== ==== 30 "​systèmes informatiques"​ et un seul programme ====
Ligne 67: Ligne 68:
 Nous allons illustrer 5 révolutions technologiques de ces 30 dernières années en exécutant (presque) un même programme sur 30 "​systèmes informatiques"​ différents : 16 d'​entre eux seront des systèmes à base de processeurs traditionnels (un ou plusieurs CPU ou Central Processing Units), 13 à base de circuits graphiques (des GPU ou GPGPU pour Graphical Processing Units ou General Purpose GPU) et le 30ème sera le seul accélérateur de type MIC (Many Integrated Cores). Nous allons illustrer 5 révolutions technologiques de ces 30 dernières années en exécutant (presque) un même programme sur 30 "​systèmes informatiques"​ différents : 16 d'​entre eux seront des systèmes à base de processeurs traditionnels (un ou plusieurs CPU ou Central Processing Units), 13 à base de circuits graphiques (des GPU ou GPGPU pour Graphical Processing Units ou General Purpose GPU) et le 30ème sera le seul accélérateur de type MIC (Many Integrated Cores).
  
-Parmi les 16 processeurs testés, 11 Intel et 5 AMD, des modèles entre 1989 et 2018. Pour les processeurs Intel : 80386SX, 80486SX, 80486DX4, Pentium 4 Northwood, E5440, X5550, E5-2637v4, E5-2680v4, Gold 5122, Silver 4144 et W-2145. Pour les AMD : Amd5x86, K6-2, K7, AthlonX2, ThreadRipper 1950X. Leurs fréquence va de 25 MHz pour le plus lent à 3.7 GHz pour le plus rapide.+Parmi les 16 processeurs testés, 11 Intel et 5 AMD, des modèles entre 1989 et 2018. Pour les processeurs Intel : 80386SX, 80486SX, 80486DX4, Pentium 4 Northwood, E5440, X5550, E5-2637v4, E5-2680v4, Gold 5122, Silver 4144 et W-2145. Pour les AMD : Amd5x86, K6-2, K7, AthlonX2, ThreadRipper 1950X. Leurs fréquences vont de 25 MHz pour le plus lent à 3.7 GHz pour le plus rapide.
  
 Pour la panoplie des GPU, nous avons 5 générations successives de circuits graphiques de gamer Nvidia de 2013 à 2019 avec les cartes GTX 560 Ti, GTX 780 Ti, GTX 980 Ti, GTX 1080 Ti, RTX Titan. Pour les circuits Nvidia graphiques "​dédiés"​ au calcul, nous avons 5 générations Tesla : C1060, M2090, K40m, P100, V100. Pour les AMD, nous avons les trois dernières générations : Nano Fiji, Vega 64 et Radeon VII. A ces 13 accélérateurs de type GPU s'​ajoute le Xeon Phi Intel 7120P. Pour la panoplie des GPU, nous avons 5 générations successives de circuits graphiques de gamer Nvidia de 2013 à 2019 avec les cartes GTX 560 Ti, GTX 780 Ti, GTX 980 Ti, GTX 1080 Ti, RTX Titan. Pour les circuits Nvidia graphiques "​dédiés"​ au calcul, nous avons 5 générations Tesla : C1060, M2090, K40m, P100, V100. Pour les AMD, nous avons les trois dernières générations : Nano Fiji, Vega 64 et Radeon VII. A ces 13 accélérateurs de type GPU s'​ajoute le Xeon Phi Intel 7120P.
Ligne 73: Ligne 74:
 ==== Echantillon de machines pertinent pour la loi de Moore ? ==== ==== Echantillon de machines pertinent pour la loi de Moore ? ====
  
-Comme premier travail ​cherchons ​d'​abord à vérifier si notre échantillon de processeurs informatiques "​respecte"​ la loi de Moore. Ce premier travail de bibliographie n'est pas si anodin ​que cela+Notre premier travail ​consiste ​d'​abord à vérifier si notre échantillon de processeurs informatiques "​respecte"​ la loi de Moore. Ce premier travail de bibliographie n'est pas si anodin ​qu'il y paraît.
  
-En effet, avant 2010, tout allait bien : d'un côté, tous les processeurs sont bien documentés et la base de données Ark d'​Intel publiait sans souci le nombre de transistors sur ses sockets de processeurs. Ainsi, nous apprenons que le E5440 dispose de 810 millions de transistors et le X5550 de 731 millions. Nous pouvons être surpris de voir le nombre de transistors baisser alors que le nombre de coeurs est identique, ​surtout ​que ce dernier embarque le contrôleur mémoire alors que son aîné ​avait cela sur son processeurCeci illustre un élément intéressant,​ lequel apparaît nettement si nous regardons une photographie de la lithographie d'un processeur : la mémoire cache "​occupe"​ une surface imposante des processeurs,​ et donc un nombre croissant de transistors : autour de 40% pour le E5440 et 30% pour le X5550. A partir de la génération suivante, Westmere, Intel est beaucoup moins prolixe sur le nombre de transistors et les informations deviennent de plus en plus difficiles à trouver avec la croissance du nombre de coeurs. Etablir une relation entre le nombre de coeurs et le nombre de transistors,​ pour une génération donnée de processeur, tient de l'ingenérie ​inverse : des informations fragmentaires sont à dénicher sur des sites spécialisés,​ en les recoupant, et, surtout, en sachant résoudre des systèmes d'​équations ! Ce sont donc des estimations pour certains d'​entre eux.+En effet, avant 2010, tout allait bien : d'un côté, tous les processeurs sont bien documentés et la base de données Ark d'​Intel publiait sans souci le nombre de transistors sur ses sockets de processeurs. Ainsi, nous apprenons que le E5440 dispose de 810 millions de transistors et le X5550 de 731 millions. Nous pouvons être surpris de voir le nombre de transistors baisser alors que le nombre de coeurs est identique, ​d'​autant plus que ce dernier embarque le contrôleur mémoire alors que son aîné ​en était dépourvuCela illustre un élément intéressant,​ lequel apparaît nettement si nous regardons une photographie de la lithographie d'un processeur : la mémoire cache "​occupe"​ une surface imposante des processeurs,​ et donc un nombre croissant de transistors : autour de 40% pour le E5440 et 30% pour le X5550. A partir de la génération suivante, Westmere, Intel est beaucoup moins prolixe sur le nombre de transistors et les informations deviennent de plus en plus difficiles à trouver avec la croissance du nombre de coeurs. Etablir une relation entre le nombre de coeurs et le nombre de transistors,​ pour une génération donnée de processeur, tient de l'ingénierie ​inverse : des informations fragmentaires sont à dénicher sur des sites spécialisés,​ en les recoupant, et, surtout, en sachant résoudre des systèmes d'​équations ! Ce sont donc des estimations pour certains d'​entre eux.
  
 Ainsi, nous avons 275000 transistors pour le 80386SX, le 80486SX est le premier à dépasser le million, le K6-2 approche les 10 millions. Le milliard n'est dépassé qu'en associant 2 E5440 Harpertown et le plus "​lourd"​ reste le ThreadRipper hébergeant près de 10 milliards de transistors. Ainsi, nous avons 275000 transistors pour le 80386SX, le 80486SX est le premier à dépasser le million, le K6-2 approche les 10 millions. Le milliard n'est dépassé qu'en associant 2 E5440 Harpertown et le plus "​lourd"​ reste le ThreadRipper hébergeant près de 10 milliards de transistors.
Ligne 81: Ligne 82:
 A la représentation de l'​évolution du nombre de transistors des "​systèmes processeurs",​ nous ajoutons leur évolution en fréquence. Nous avons ajouté la "​fréquence turbo"​. Cette fréquence est très significative pour les processeurs récents, représentant jusqu'​à 40% de la fréquence nominale. Elle est essentiellement exploitée pour les tâches non parallélisées. A la représentation de l'​évolution du nombre de transistors des "​systèmes processeurs",​ nous ajoutons leur évolution en fréquence. Nous avons ajouté la "​fréquence turbo"​. Cette fréquence est très significative pour les processeurs récents, représentant jusqu'​à 40% de la fréquence nominale. Elle est essentiellement exploitée pour les tâches non parallélisées.
  
-Une représentation logarithmique permet de saisir la progression géométrique du nombre de transistors. Pour la fréquence, elle suit aussi une progression géométrique mais, en 2005, s'​arrête,​ régresse et repart légèrement à la hausse. Avec une méthode de régression de type moindre ​carrés, tâchons de retrouver les "​lois"​ qui gouvernent ces droites.+Une représentation logarithmique permet de saisir la progression géométrique du nombre de transistors. Pour la fréquence, elle suit aussi une progression géométrique mais, en 2005, elle s'​arrête, ​puis régresse et repart légèrement à la hausse. Avec une méthode de régression de type "​moindres ​carrés", tâchons de retrouver les "​lois"​ qui gouvernent ces droites.
  
 {{ :​developpement:​activites:​qualification:​transistorsfrequencies.png?​500 |}} {{ :​developpement:​activites:​qualification:​transistorsfrequencies.png?​500 |}}
Ligne 87: Ligne 88:
 Pour le nombre de transistors,​ si nous exprimons cette loi par ''​Transistors(Y)=exp(a*Y+b)'',​ nous trouvons que ''​a''​ vaut ''​0.342''​ et ''​b''​ vaut ''​-667''​. Si nous cherchons le temps qui sépare le doublement du nombre de transistors par cette loi, nous effectuons le calcul ''​ln(2)/​a''​ et nous trouvons ''​2.02'',​ soit 2 à 1% près ! Effectivement,​ nous retrouvons la loi selon laquelle l'​intégration de transistors sur un support double tous les deux ans : la loi de Moore. Petit détail comique : si nous cherchons quand le nombre de transistors était égal à l'​unité (par ''​-b/​a''​),​ nous tombons sur 1951 (seulement 3 ans après l'​invention du transistors :-). Pour le nombre de transistors,​ si nous exprimons cette loi par ''​Transistors(Y)=exp(a*Y+b)'',​ nous trouvons que ''​a''​ vaut ''​0.342''​ et ''​b''​ vaut ''​-667''​. Si nous cherchons le temps qui sépare le doublement du nombre de transistors par cette loi, nous effectuons le calcul ''​ln(2)/​a''​ et nous trouvons ''​2.02'',​ soit 2 à 1% près ! Effectivement,​ nous retrouvons la loi selon laquelle l'​intégration de transistors sur un support double tous les deux ans : la loi de Moore. Petit détail comique : si nous cherchons quand le nombre de transistors était égal à l'​unité (par ''​-b/​a''​),​ nous tombons sur 1951 (seulement 3 ans après l'​invention du transistors :-).
  
-Quant est-il de l'​évolution de la fréquence ? Nous constatons 2 périodes successives : une première suit une loi à peu près comparable de 1989 à 2004. En exprimant cette loi par un ''​Frequence(Y)=exp(c*Y+d)'',​ nous trouvons ''​c=0.307''​ et ''​d=-609''​. La fréquence a donc doublé tous les 27 mois en moyenne durant cette période. ​Après, nous assistons à un léger ​replis ​sous les 3 GHz puis une réaugmentation ces toutes dernières années. Cette "​brisure"​ est fondamentale : elle illustre comment l'​effet Joule a stoppé net la progression de la fréquence alors que la miniaturisation des transistors se poursuivait. Les transistors devaient donc être exploité ​"​autrement"​.+Qu'​en ​est-il de l'​évolution de la fréquence ? Nous constatons 2 périodes successives : une première suit une loi à peu près comparable de 1989 à 2004. En exprimant cette loi par un ''​Frequence(Y)=exp(c*Y+d)'',​ nous trouvons ''​c=0.307''​ et ''​d=-609''​. La fréquence a donc doublé tous les 27 mois en moyenne durant cette période. ​Ensuite, nous assistons à un léger ​repli sous les 3 GHz puis une réaugmentation ces toutes dernières années. Cette "​brisure"​ est fondamentale : elle illustre comment l'​effet Joule a stoppé net la progression de la fréquencealors que la miniaturisation des transistors se poursuivait. Les transistors devaient donc être exploités ​"​autrement"​.
  
 Toujours est-il que cette simple représentation montre que notre échantillon est bien pertinent : il respecte la loi de Moore et nous montre cette "​transition"​ en 2004 sur la fréquence. Reste maintenant à exécuter notre code. Toujours est-il que cette simple représentation montre que notre échantillon est bien pertinent : il respecte la loi de Moore et nous montre cette "​transition"​ en 2004 sur la fréquence. Reste maintenant à exécuter notre code.
Ligne 93: Ligne 94:
 ==== Quelles révolutions informatiques ces 30 dernières années ? ==== ==== Quelles révolutions informatiques ces 30 dernières années ? ====
  
-Il ne suffit pas de disposer ​de plus de transistors sur un processeur pour augmenter sa capacité de traitement. Cependant, cette croissance complémentaire va permettre d'​intégrer directement dans le circuit des éléments qui ne s'y trouvaient pas, ou seulement dans des architectures "​confidentielles"​ pour l'​époque. Voici les 5 (r)évolutions que nous avons choisies :+Il ne suffit pas de disposer de transistors ​en plus sur un processeur pour augmenter sa capacité de traitement. Cependant, cette croissance complémentaire va permettre d'​intégrer directement dans le circuit des éléments qui ne s'y trouvaient pas, ou seulement dans des architectures "​confidentielles"​ pour l'​époque. Voici les 5 (r)évolutions que nous avons choisies :
   * intégration systématique d'une unité de calcul en virgule flottante, la FPU (ou Floating Point Unit)   * intégration systématique d'une unité de calcul en virgule flottante, la FPU (ou Floating Point Unit)
   * exploitation de l'​architecture RISC en interne au processeur   * exploitation de l'​architecture RISC en interne au processeur
Ligne 101: Ligne 102:
  
 Chacune de ces révolutions sera illustrée par un avant et un après : Chacune de ces révolutions sera illustrée par un avant et un après :
-  * Intégration FPU. Avant : le 30386SX et le 80486SX. Après le 80486DX4 et le Amd5x86 +  * Intégration FPU. Avant : le 30386SX et le 80486SX. Après ​le 80486DX4 et le Amd5x86 ​; 
-  * Exploitation RISC et vectorisation. Avant : les précédents. Après : le K6-2 +  * Exploitation RISC et vectorisation. Avant : les précédents. Après : le K6-2 ; 
-  * Multiplication des coeurs. Avant : les précédents. Après : le Northwood P4 et le AthlonX2+  * Multiplication des coeurs. Avant : les précédents. Après : le Northwood P4 et le AthlonX2 ​;
   * Manipulation des GPU. Avant : tous les processeurs. Après : 10 ans d'​évolution de GPU et GPGPU.   * Manipulation des GPU. Avant : tous les processeurs. Après : 10 ans d'​évolution de GPU et GPGPU.
- 
- 
  
 === Conditions expérimentales et exécution du code === === Conditions expérimentales et exécution du code ===
  
-C'est un même code que nous allons exploiter mais ce dernier propose un certain nombre d'​options. ​Parmi ces options, nous avons +C'est un même code que nous allons exploiter mais ce dernier propose un certain nombre d'​options. ​Options parmi lesquelles, nous avons : 
-  * la masse (en fait, un trou noir statique n'est défini "​que"​ par ce paramètre) +  * la masse (en fait, un trou noir statique n'est défini "​que"​ par ce paramètre) ​; 
-  * la taille de l'​image générée +  * la taille de l'​image générée ​; 
-  * les rayons intérieur et extérieur du disque de plasma+  * les rayons intérieur et extérieur du disque de plasma ​;
   * l'​inclinaison de l'​observateur par rapport au disque   * l'​inclinaison de l'​observateur par rapport au disque
   * la nature de l'​émission du plasma :   * la nature de l'​émission du plasma :
-    * émission monochromatique,​ noté Mono, irréaliste mais instructive +    * émission monochromatique,​ noté Mono, irréaliste mais instructive, 
-    * émission de corps noir, noté BB, plus réaliste+    * émission de corps noir, noté BB, plus réaliste.
  
 En fonction de la performance et de la mémoire disponible dans chaque système. En fonction de la performance et de la mémoire disponible dans chaque système.
 +
 +=== Eléments de sortie ===
 +
 +Ces simulations ont pour objectif de "​montrer"​ quelle serait l'​apparence d'un disque d'​accrétion gravitant autour d'un trou noir. Cependant, ce n'est pas une mais deux images que nous allons synthétiser :
 +  * la première en **[https://​fr.wikipedia.org/​wiki/​D%C3%A9calage_vers_le_rouge|z]]**
 +  * la seconde en **[[https://​fr.wikipedia.org/​wiki/​Flux_lumineux|flux]]**
 +
 +Si le programme originel ne "​produit"​ qu'une image en niveaux de gris normalisés sur la valeur maximale, le programme en Python exploite des "​fausses couleurs"​ de manière équivalente à celle produite en avril dernier.
 +
 +Voici une image en **décalage spectral** en fausses couleurs, pour un observateur situé 10° au dessus du disque.
 +
 +{{ :​developpement:​activites:​qualification:​trounoirz_trajectopixel_device0_casimir_20191217_104144.png?​500 |}}
 +
 +Voici une image en **flux** dont l'​émission du disque est une **raie monochromatique** (une unique couleur).
 +
 +{{ :​developpement:​activites:​qualification:​trounoirf_trajectopixel_device0_casimir_20191217_104144.png?​500 |}}
 +
 +Voici une image en **flux** dont l'​émission est un **spectre de corps noir** (une distribution de couleurs).
 +
 +{{ :​developpement:​activites:​qualification:​trounoirf_trajectopixel_device0_casimir_20191217_104150.png?​500 |}} 
  
 === Intégration du calcul flottant dans le processeur === === Intégration du calcul flottant dans le processeur ===
  
-Dès le premier IBM PC, Intel a proposé sur sa carte mère d'​associer au processeur une unité spécialisée permettant d'​effectuer les opérations classiques de calcul en virgule flottante : le [[https://​en.wikipedia.org/​wiki/​Intel_8087|8087]] était ainsi associé au processeur [[https://​en.wikipedia.org/​wiki/​Intel_8086|8086]]. Cette approche s'est poursuivie durant toutes les années 1980 sur les évolutions de ses processeurs,​ le 80286 puis le 80386. Il suffisait d'​ajouter un x87 (80287 ou 80387) pour "​transformer des minutes en secondes"​.+Dès le premier IBM PC, Intel a proposé sur sa carte-mère d'​associer au processeur une unité spécialisée permettant d'​effectuer les opérations classiques de calcul en virgule flottante : le [[https://​en.wikipedia.org/​wiki/​Intel_8087|8087]] était ainsi associé au processeur [[https://​en.wikipedia.org/​wiki/​Intel_8086|8086]]. Cette approche s'est poursuivie durant toutes les années 1980 sur les évolutions de ses processeurs,​ le 80286 puis le 80386. Il suffisait d'​ajouter un x87 (80287 ou 80387) pour "​transformer des minutes en secondes"​.
  
 Dans notre cas, nous allons exploiter 4 processeurs : deux sont dépourvus de coprocesseur flottant (le 80386SX à 40 MHz et le 80486SX à 25 MHz), deux en sont équipés (le 80486DX à 75 MHz et le Amd5x86 à 133 MHz). Dans notre cas, nous allons exploiter 4 processeurs : deux sont dépourvus de coprocesseur flottant (le 80386SX à 40 MHz et le 80486SX à 25 MHz), deux en sont équipés (le 80486DX à 75 MHz et le Amd5x86 à 133 MHz).
Ligne 131: Ligne 150:
 {{:​developpement:​activites:​qualification:​8038680486.png?​500|}} {{:​developpement:​activites:​qualification:​8038680486.png?​500|}}
  
-Nous constatons que la fréquence ne fait pas tout ! En effet, le 80386SX est à une fréquence presque 2 fois supérieure au 80486SX et pourtant le 80486SX l'​emporte largement. ​Plus de transistors ont permis au 80486SX de multiplier par 2.5x la performance dans les deux cas. +Nous constatons que la fréquence ne fait pas tout ! En effet, le 80386SX est à une fréquence presque 2 fois supérieure au 80486SX et pourtant le 80486SX l'​emporte largement. ​D'​avantage ​de transistors ont permis au 80486SX de multiplier par 2.5x la performance dans les deux cas. 
  
 Si nous ajoutons d'​abord le premier processeur de notre gamme à être équipé de FPU, le 80486DX4, le gain est très significatif : Si nous ajoutons d'​abord le premier processeur de notre gamme à être équipé de FPU, le 80486DX4, le gain est très significatif :
Ligne 225: Ligne 244:
  
 Ce constat est sans appel : il va falloir exploiter tous les coeurs disponibles,​ logiques ou physiques, pour augmenter la performance de son code. Ce constat est sans appel : il va falloir exploiter tous les coeurs disponibles,​ logiques ou physiques, pour augmenter la performance de son code.
 +
 +===== 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.
  
 ==== Parallélisation à la OpenMP ==== ==== Parallélisation à la OpenMP ====
Ligne 240: Ligne 263:
 Il est alors intéressant,​ avec les commandes "​système",​ de contrôler si tous les coeurs sont bien exploités efficacement. Il est alors 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.
 +
 +{{ :​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.
 +
 +Regardons s'il en est de même en Mono.
 +
 +{{ :​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.
 +
 +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.
 +
 +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.
 +
 +{{ :​developpement:​activites:​qualification:​tr_omp_bb.png?​500 |}} 
 +
 +La courbe violette montre la performance pour le code non parallélisé. Nous distinguons que pour une parallélisation fixée au nombre de coeurs physiques, la performance ne triple même pas alors que nous avons 16 coeurs. Si nous doublons le nombre de tâches concurrentes,​ la performance approche seulement 5x celle du code séquentiel. Par contre, pour une parallélisation imposée sur 64x le nombre de coeurs, la performance dépasse 16x.
 +
 +Regardons s'il en est de même pour ma méthode Mono.
 +
 +{{ :​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.
 +
 +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.
 +
 +==== 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. ​
 +
 +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).
 +
 +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.
 +
 +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 !
 +
 +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é.
 +
 +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 :
 +  * **Original** : 1
 +  * **EachCircle** : N/2
 +  * **EachPixel** : NxN
 +  * **TrajectoCircle** : N/2 puis 8xNxN
 +  * **TrajectoPixel** : N/2 puis NxN
 +
 +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**). ​
 +
 +=== 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.
 +
 +{{ :​developpement:​activites:​qualification:​tr1950x_7methods_elapsed.png?​500 |}}
 +
 +Sur cette illustration,​ //less is better//. Nous constatons que la méthode **OpenCL EachPixel** est fortement inefficace : le calcul élémentaire de la trajectoire sur chaque pixel est bien rédhibitoire. De plus, la méthode **OpenCL Original** est moins efficace qu'​avec le programme séquentiel **Serial**. Les méthodes hybrides **OpenCL TrajectoPixel** et **OpenCL TrajectoCircle** sont plutôt proches. Les méthodes **OpenMP** et **OpenCL EachCircle**,​ comparables en approche de programmation,​ semblent tourner à l'​avantage de cette dernière. ​
 +
 +De manière à en avoir le coeur net, calculons la performance,​ le //​PixHertz//​ de chaque méthode en divisant le nombre de pixels de l'​image par les deux durées **Compute** et **Elapsed**.
 +
 +{{ :​developpement:​activites:​qualification:​tr1950x_7methods_pixhertz.png?​500 |}}
 +
 +Nous confirmons que la méthode **OpenCL EachPixel** est bien la plus mauvaise, que **OpenCL Original** est un peu moins bon que **Serial**. Sur la comparaison entre **OpenMP** et **OpenCL EachCircle**,​ nous remarquons que OpenCL l'​emporte de presque un facteur 2, mais uniquement en calcul. En temps écoulé, **EachCircle** reste comparable à **OpenMP**... En effet, le temps de transfert des plusieurs GB qu'​occupent les deux images de plus d'1/4 de milliard de pixels prend autour d'une seconde, à peine supérieur au temps de calcul : une belle illustration que, en OpenCL, le temps de transfert entre hôte et périphérique "​compte"​ pour la performance du calcul. Nous sommes toutefois rassurés que notre effort de portage OpenCL nous permette, en calcul, d'​améliorer nos performances.
 +
 +=== Les performances sur tous les processeurs en OpenCL ===
 +
 +Focalisons nous sur la capacité de calcul, donc le "​Compute Time" pour les simulations **Mono** ou **BB**.
 +
 +{{ :​developpement:​activites:​qualification:​cpu_compute_mono.png?​500 |}}
 +
 +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 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.
 +
 +Regardons ces performances maintenant en intégrant les temps de transferts
 +
 +{{ :​developpement:​activites:​qualification:​cpu_elapsed_mono.png?​500 |}}
 +
 +En Mono, le Threadripper se retrouve au même niveau en OpenCL qu'en OpenMP. Seul le dernier Skylake Intel W-2145 gagne un facteur 2.5x entre OpenCL et OpenMP. Les autres gagnent timidement de 80% à quelques dizaines de %. Le Silver déçoit en étant moins performant en OpenCL qu'en OpenMP.
 +
 +{{ :​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 ​
 +
 +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.
 +
 +=== Intégration des accélérateurs en OpenCL ===
 +
 +== Première comparaison avec la première Tesla ==
 +
 +En 2008, Nvidia titre [[https://​www.nvidia.com/​docs/​IO/​43395/​NV_DS_Tesla_PSC_US_Mar09_LowRes.pdf|Tesla Personal Supercomputer]]. Dans une machine au format "​station de travail",​ de 3 à 4 Nvidia Tesla C1060. Avec un puissance estimée à 4 TFlops, l'​article précise que cette station équivaut à 250 ordinateurs personnels.
 +
 +La première étape consiste donc à comparer cette machine avec une station de travail gonflée ou un serveur bisocket équivalent. A l'​époque,​ le processeur "fer de lance" de Intel est le Harpertown, un assemblage de 2 Core2 Duo Penryn.
 +
 +{{ :​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.
 +
 +== AMD : les trois dernières générations ==
 +
 +Nous nous propulsons en 2019. Face aux meilleurs de nos processeurs (le Threadripper 1950X et l'​Intel C4145), nous opposons les trois dernières générations AMD : Nano Fury, Vega 64 et Radeon 7 (ou Vega 2).
 +
 +{{ :​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.
 +
 +== 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).
 +
 +{{ :​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.
 +
 +== Et les accélérateurs non GPU dans tout ça ? ==
 +
 +Nous avons évoqué le Xeon Phi, le coprocesseur arithmétique d'​Intel disposant de 60 coeurs de calcul. Comment le placer dans ce classement ? La solution est d'​abord de le replacer dans son contexte, en 2013-2014 avec des processeurs et des GPU comparables : une machine bi-sockets avec processeurs Westmere et les GPU ou GPGPU GTX780Ti et Tesla K40m.
 +
 +{{ :​developpement:​activites:​qualification:​xeonphivsall.png?​500 |}}
 +
 +Nous constatons que le Xeon Phi 7120P est à peine meilleur en Mono que le système bi-sockets de son époque, mais 2x plus lent que la GTX ou la Tesla. Par contre, il est infiniment meilleur en BB, dépassant même la GTX. Ainsi, le ratio de performances est l'un des meilleurs que nous ayons testés. C'est cependant une pâle satisfaction pour la carte qui devait révolutionner l'​informatique.
 +
 +== Tesla face aux autres... ==
 +
 +Nous avons brièvement testé les Tesla C1060 et Tesla K40m en les comparant aux accélérateurs et processeurs de leur époque. Il est temps de comparer ces 5 générations de Tesla, C1060 (circuit GT200), M2090 (circuit Fermi), K40m (circuit Kepler), P100 (circuit Pascal), V100 (circuit Volta) aux autres GPU de gamer ou AMD.
 +
 +{{ :​developpement:​activites:​qualification:​voltavsall.png?​500 |}}
 +
 +Nous constatons que les Tesla offrent des performances comparables aux cartes de gamer en Mono. La très onéreuse Tesla V100 est même inférieure à l'​onéreuse RTX Titan. Il en est de même de la Tesla P100 avec sa concurrente directe, la GTX 1080 Ti. Par contre, dès que nous passons en BB, le rapport s'​inverse complètement : les Tesla récentes sont entre 5x et 10x plus rapides que les GTX ou RTX équivalentes. ​
 +
 +D'où vient une telle différence ? Il y a la charge calculatoire d'​abord : beaucoup plus d'​opérations sont nécessaires dans la simulation BB face à Mono. Il y a ensuite la nature des opérations : beaucoup de fonctions transcendantes sont exploitées et, manifestement,​ elles ne sont pas confiées aux mêmes unités de calculs. Une telle différence suggère que les opérations transcendantes sont confiées systématiquement aux unités flottantes 64 bits beaucoup moins nombreuses sur les cartes de Gamer.
 +
 +D'​autre part, un espoir du côté de AMD : jusqu'​alors complètement distancée dans toutes les simulations,​ la Radeon 7 dépasse toutes les cartes de gamer Nvidia, RTX Titan compris, et ce d'un facteur 2, mais uniquement en BB.
 +
 +=== Et CUDA dans tout cela ? ===
 +
 +Pour l'​instant,​ nous nous sommes contentés d'​exploiter le même code gépufié pour tous nos "​périphériques"​ OpenCL, qu'ils soient processeurs,​ GPU, GPGPU de marque Nvidia ou AMD, voire l'​accélérateur Intel.
 +
 +Est-il possible d'​améliorer ses performances,​ celles de Nvidia, en exploitant le langage originel créé par Nvidia pour "​parler"​ au GPU ? La réponse est oui, évidemment,​ mais cela exige un petit portage. Si la nature des "​noyaux"​ de calcul est tout à fait comparable (c'est du C plutôt très simple), la manière de distriubuer le travail n'est pas spécialement comparable.
 +
 +== Un premier jet simple, mais peu convaincant ==
 +
 +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.
 +
 +== Exploiter les deux étages de parallélisme : une nécessité mais comment ? ==
 +
 +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 :
 +<​code>​
 +uint xi=(uint)get_global_id(0);​
 +</​code>​
 +Alors que nous avons en CUDA :
 +<​code>​
 +uint xi=(uint)(blockIdx.x*blockDim.x+threadIdx.x);​
 +</​code>​
 +
 +Pour le même test précédent sur la Tesla P100, en distribuant sur 32 //threads// et 32 par 32 (soit 1024) //threads// les deux tâches, nous sommes, en CUDA, légèrement meilleur que OpenCL sur une simulation BB (20% plus rapide). Par contre, le gain est énorme en Mono puisque nous passons d'une seconde à 0.06 seconde, soit 13 fois plus rapide.
 +
 +Reste maintenant à étendre cela à tous nos GPU et GPGPU Nvidia !
 +
 +== 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.
 +
 +{{ :​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.
 +
 +{{ :​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.
 +
 +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.
 +
 +=== 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 !
 +
 +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.
 +
 +{{ :​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.
 +
 +{{ :​developpement:​activites:​qualification:​cudavsall_elapsed_mono.png?​500 |}}
 +
 +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).
 +
 +==== Comparaison finale sur 30 ans de matériel informatique ====
 +
 +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.
 +
 +{{ :​developpement:​activites:​qualification:​performancesyear_bb.png?​600 |}}
 +
 +Pour les simulations BB, Les triangles jaunes (programme OpenCL) se placent bien au dessus des carrés bleus (programme séquentiel) voire des losanges oranges (programme OpenMP). L'​échelle étant logarithmique en ordonnée, nous voyons également émerger timidement quelques triangles verts (programme CUDA).
 +
 +{{ :​developpement:​activites:​qualification:​performancesyear_mono.png?​600 |}}
 +
 +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.txt · Dernière modification: 2019/12/17 15:34 par equemene