developpement:activites:qualification:30ans1code

Différences

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

--- developpement:activites:qualification:30ans1code [2019/12/16 12:09]
equemene [Parallélisation avec OpenCL]
+++ developpement:activites:qualification:30ans1code [2019/12/16 14:28]
equemene [Parallélisation avec OpenCL]
@@ Ligne 390: / Ligne 390: @@
 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 u peu supérieur.
+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 !
-En effet, pour exploiter efficacement CUDA, il faut IMPERATIVEMENT solliciter les "Threads".

developpement/activites/qualification/30ans1code.txt · Dernière modification: 2019/12/17 15:34 par equemene

Rechercher

Translations

Navigation

Piste:

Boîte à outils