developpement:activites:qualification:30ans1code

Différences

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

--- developpement:activites:qualification:30ans1code [2019/11/27 17:36]
equemene [Parallélisation avec OpenCL]
+++ developpement:activites:qualification:30ans1code [2019/12/17 10:52]
equemene
@@ Ligne 14: / Ligne 14: @@
 (détournement de Tolkien...)
 ===== Résumé =====
@@ Ligne 105: / Ligne 106: @@
   * 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.
 === Conditions expérimentales et exécution du code ===
@@ Ligne 120: / Ligne 119: @@
 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 ===
@@ Ligne 317: / Ligne 336: @@
 === 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*.
+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 |}}
@@ Ligne 329: / Ligne 348: @@
 Regardons ces performances maintenant en intégrant les temps de transferts
-{{ :developpement:activites:qualification:cpu_compute_bb.png?500 |}}
+{{ :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 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
+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

Rechercher

Translations

Navigation

Piste:

Boîte à outils