Ci-dessous, les différences entre deux révisions de la page.
Les deux révisions précédentes Révision précédente Prochaine révision | Révision précédente Dernière révision Les deux révisions suivantes | ||
formation:gpu4cbp [2023/03/08 10:44] equemene [Retour au C et ses implémentations OpenMP et OpenACC] |
formation:gpu4cbp [2023/03/13 08:58] equemene |
||
---|---|---|---|
Ligne 700: | Ligne 700: | ||
Pour cela, reprenons notre programme de simple addition ''MySteps_1.py''. Il dispose de son implémentation C dans le dossier ETSN au même titre que tous les autres ''MySteps_1.c''. Cette implémentation C reprend tous les éléments et permet de "juger" de la vitesse d'exécution sur une langage compilé. L'unique paramètre à l'appel du programme est la taille des vecteurs. Pour la compilation de ce programme :<code>gcc -O3 -o MySteps_1 MySteps_1.c -lm</code> | Pour cela, reprenons notre programme de simple addition ''MySteps_1.py''. Il dispose de son implémentation C dans le dossier ETSN au même titre que tous les autres ''MySteps_1.c''. Cette implémentation C reprend tous les éléments et permet de "juger" de la vitesse d'exécution sur une langage compilé. L'unique paramètre à l'appel du programme est la taille des vecteurs. Pour la compilation de ce programme :<code>gcc -O3 -o MySteps_1 MySteps_1.c -lm</code> | ||
- | Le principe de base de la programmation OpenMP consiste à "casser des boucles", où plutôt distribuer les calculs indépendants sur les éléments d'un tableau aux ressources à disposition. Pour cela, OpenMP s'appuie sur le "balisage" du code source : les ''#pragma''. Ces messages "aident" le préprocesseur avant le compilateur à modifier le code pour distribuer les tâches sur les ressources à disposition. La version OpenMP de ''MySteps_1.c'' est ''MySteps_1_openmp.c'' : il est possible par une simple commande ''diff'' entre les deux codes sources de juger des différences : un unique ''#pragma omp parallel for'' distribue les éléments de la boucle aux ressources disponibles. La compilation s'effectue en précisant l'usage de OpenMP : <code>gcc -fopenmp -O3 -o MySteps_1_openmp MySteps_1_openmp.c -lm -lgomp</code>. | + | Le principe de base de la programmation OpenMP consiste à "casser des boucles", où plutôt distribuer les calculs indépendants sur les éléments d'un tableau aux ressources à disposition. Pour cela, OpenMP s'appuie sur le "balisage" du code source : les ''#pragma''. Ces messages "aident" le préprocesseur avant le compilateur à modifier le code pour distribuer les tâches sur les ressources à disposition. La version OpenMP de ''MySteps_1.c'' est ''MySteps_1_openmp.c'' : il est possible par une simple commande ''diff'' entre les deux codes sources de juger des différences : un unique ''#pragma omp parallel for'' distribue les éléments de la boucle aux ressources disponibles. La compilation s'effectue en précisant l'usage de OpenMP : <code>gcc -fopenmp -O3 -o MySteps_1_openmp MySteps_1_openmp.c -lm -lgomp</code> |
Il est alors possible d'effectuer une comparaison de performances entre la version sérielle et la version parallélisée avec OpenMP. | Il est alors possible d'effectuer une comparaison de performances entre la version sérielle et la version parallélisée avec OpenMP. | ||
Ligne 708: | Ligne 708: | ||
Le principe de la programmation OpenACC est très proche de OpenMP : l'exploitation de balises permettant d'identifier dans le code source les calculs à envoyer sur le périphérique externe (GPU ou accélérateur). Là, les balises ne servent pas uniquement à identifier les boucles mais aussi les portions de code qui seront à "pousser" dans le périphérique externe ainsi que les éléments de tableau. La version OpenACC de ''MySteps_1.c'' est ''MySteps_1_openacc.c'' : il est possible par une simple commande ''diff'' entre les deux codes sources de juger des différences : le nombre de ''#pragma'' est plus important. Le premier, ''#pragma acc data copyin(a[0:size],b[0:size]),copyout(res[0:size])'', indique les données à "copier sur" et à "récupérer de" du périphérique externe. Le second, ''#pragma acc parallel loop'', distribue les éléments de la boucle au périphérique externe. La compilation est par contre plus "verbeuse". Elle intègre l'appel à OpenACC puis des options de compilations exploitant l'outil de compilation Nvidia ''nvcc'' :<code>gcc -O3 -fopenacc -foffload=nvptx-none -foffload="-O3 -misa=sm_35 -lm" -o MySteps_1_openacc MySteps_1_openacc.c -lm</code> | Le principe de la programmation OpenACC est très proche de OpenMP : l'exploitation de balises permettant d'identifier dans le code source les calculs à envoyer sur le périphérique externe (GPU ou accélérateur). Là, les balises ne servent pas uniquement à identifier les boucles mais aussi les portions de code qui seront à "pousser" dans le périphérique externe ainsi que les éléments de tableau. La version OpenACC de ''MySteps_1.c'' est ''MySteps_1_openacc.c'' : il est possible par une simple commande ''diff'' entre les deux codes sources de juger des différences : le nombre de ''#pragma'' est plus important. Le premier, ''#pragma acc data copyin(a[0:size],b[0:size]),copyout(res[0:size])'', indique les données à "copier sur" et à "récupérer de" du périphérique externe. Le second, ''#pragma acc parallel loop'', distribue les éléments de la boucle au périphérique externe. La compilation est par contre plus "verbeuse". Elle intègre l'appel à OpenACC puis des options de compilations exploitant l'outil de compilation Nvidia ''nvcc'' :<code>gcc -O3 -fopenacc -foffload=nvptx-none -foffload="-O3 -misa=sm_35 -lm" -o MySteps_1_openacc MySteps_1_openacc.c -lm</code> | ||
- | Une option, comme pour le cas de l'implémentation OpenMP, permete d'inhiber l'exécution sérielle :<code>gcc -DNOSERIAL -O3 -fopenacc -foffload=nvptx-none -foffload="-O3 -misa=sm_35 -lm" -o MySteps_1_openacc_NoSerial MySteps_1_openacc.c -lm</code>. | + | Une option, comme pour le cas de l'implémentation OpenMP, permete d'inhiber l'exécution sérielle :<code>gcc -DNOSERIAL -O3 -fopenacc -foffload=nvptx-none -foffload="-O3 -misa=sm_35 -lm" -o MySteps_1_openacc_NoSerial MySteps_1_openacc.c -lm</code> |
- | La question légitime est maintenant de "juger" de la différence de performance entre OpenMP et OpenACC, en langage C, face à OpenCL en Python sur CPU et GPU. | + | La question légitime est maintenant de "juger" de la différence de performance entre OpenMP et OpenACC, en langage C, face à OpenCL en Python sur CPU et GPU. Sur la machine **gtxtitan**, nous disposons d'une GTX Titan et d'une CPU E6-2620 avec 6 coeurs à 2 GHz. |
^ Size ^ C/Serial ^ C/OpenMP ^ C/OpenACC ^ Numpy ^ PyCL CPU ^ PyCL GPU ^ | ^ Size ^ C/Serial ^ C/OpenMP ^ C/OpenACC ^ Numpy ^ PyCL CPU ^ PyCL GPU ^ | ||
Ligne 735: | Ligne 735: | ||
^ 1073741824| 954322688| 1854957056| | 637516695| | | | ^ 1073741824| 954322688| 1854957056| | 637516695| | | | ||
+ | Les enseignements de ces comparaisons entre implémentations de la simple addition de deux vecteurs sont les suivantes : | ||
+ | - la C/OpenMP est meilleur mais lorsque la taille du vecteur est supérieure à 1 million | ||
+ | - la C sérielle reste la plus efficace pour les petites tailles (inférieure à 32768) | ||
+ | - la Python/Numpy reste très compétitive pour des taille entre 32768 et 1048576 | ||
+ | - la OpenACC reste supérieure aux implémentations OpenCL/CPU et OpenCL/GPU d'un facteur 3 | ||
+ | |||
+ | Se contenter uniquement de ce test inviterait à fuire Python/OpenCL. Cependant, nous avons vu dans sur ''MySteps_2.py'' que la charge calculatoire doit être "vraiment" significative pour que le Python/OpenCL l'emporte de manière significative. Nous reviendrons donc dans la suite sur des versions modifiées de ces programmes C intégrant la fonction de Mylq ''MySillyFunction'', appelée plusieurs fois, pour juger si Python/OpenCL reste compétitif face à OpenMP et OpenACC. | ||
===== Un intermède CUDA et son implémentation PyCUDA ===== | ===== Un intermède CUDA et son implémentation PyCUDA ===== | ||
Ligne 856: | Ligne 863: | ||
* ''MySteps_5b.py'' : intègre une utilisation hybride des //Blocks// et des //Threads// | * ''MySteps_5b.py'' : intègre une utilisation hybride des //Blocks// et des //Threads// | ||
* ''MySteps_5c.py'' : augmente la complexité arithmétique d'un facteur 16 | * ''MySteps_5c.py'' : augmente la complexité arithmétique d'un facteur 16 | ||
+ | |||
+ | ===== Comparaison de toutes les implémentations : victoire incontestée de OpenCL ===== | ||
+ | |||
+ | Nous indiquions dans la "timide" implémentation de ''MySteps_1'' en C, OpenMP/C et OpenACC/C que nous allions revenir pour comparer finalement les 6 implémentations d'un même algorithme "empâté" d'addition de vecteurs, histoire de véritablement fixer les idées. De manière à juger l'influence de la charge, nous avons intégré le recours successifs à la //fonction de Mylq// ''MySillyFunction'' pour moduler la "charge calculatoire" de chaque élément de vecteur de l'addition. | ||
+ | |||
+ | Ces 6 implémentations sont réparties dans 4 programmes : | ||
+ | * ''MySteps_6.c'' : implémentation sérielle en C | ||
+ | * ''MySteps_6_openmp.c'' : implémentation OpenMP en C | ||
+ | * ''MySteps_6_openacc.c'' : implémentation OpenACC en C | ||
+ | * ''MySteps_6.py'' : implémentations Numpy, OpenCL et CUDA en Python | ||
+ | |||
+ | Les trois implémentations en C ont comme argument la taille et le nombre d'applications de la fonction de Mylq. | ||
+ | |||
+ | Les trois implémentations en Python ''MySteps_6.py'' dans un même programme disposent des arguments suivants : | ||
+ | * ''-h'' la documentation sommaire avec la liste des périphériques | ||
+ | * ''-s'' la taille du vecteur | ||
+ | * ''-d'' le périphérique, à choisir entre ''0'' et ''4'' généralement | ||
+ | * ''-g'' le choix de l'implémentation ''OpenCL'' ou ''CUDA'' | ||
+ | * ''-c'' le nombre d'appels successifs de la fonction de Mylq | ||
+ | * ''-t'' le nombre de //Threads// en CUDA | ||
+ | * ''-n'' la suspension de l'exécution en "mode natif" | ||
+ | |||
+ | Ainsi, pour une exécution sur une machine disposant d'un GPU avec 6GB de RAM, l'exécution de <code>./MySteps_6.py -d 0 -g OpenCL -s $((2**25)) -c 10</code> donne :<code> | ||
+ | Device Selection : 0 | ||
+ | GpuStyle used : OpenCL | ||
+ | Size of complex vector : 33554432 | ||
+ | Number of silly calls : 10 | ||
+ | Number of Threads : 1 | ||
+ | Serial compute : 1 | ||
+ | Device #0 from NVIDIA Corporation of type xPU : NVIDIA GeForce GTX TITAN | ||
+ | Device #1 from NVIDIA Corporation of type xPU : Quadro K420 | ||
+ | Device #2 from The pocl project of type xPU : pthread-Intel(R) Xeon(R) CPU E5-2620 0 @ 2.00GHz | ||
+ | Device #3 from Advanced Micro Devices, Inc. of type xPU : Intel(R) Xeon(R) CPU E5-2620 0 @ 2.00GHz | ||
+ | Device #4 from Intel(R) Corporation of type xPU : Intel(R) Xeon(R) CPU E5-2620 0 @ 2.00GHz | ||
+ | NativeRate: 128342 | ||
+ | CPU/GPU selected: NVIDIA GeForce GTX TITAN | ||
+ | Copy from Host 2 Device : 0.147 | ||
+ | Building kernels : 0.150 | ||
+ | Allocation on Device for results : 0.000 | ||
+ | Synthesis of kernel : 0.014 | ||
+ | Execution of kernel : 0.321 | ||
+ | Allocation on Host for results: 0.000 | ||
+ | Copy from Device 2 Host : 0.073 | ||
+ | OpenCLRate: 34904715 | ||
+ | [ 1.1920929e-07 0.0000000e+00 0.0000000e+00 ... -1.1920929e-07 | ||
+ | 0.0000000e+00 1.1920929e-07] | ||
+ | 0.016262252 | ||
+ | Results between Native & OpenCL seem to be too different! | ||
+ | OpenCLvsNative ratio: 271.966426 | ||
+ | </code> | ||
+ | |||
+ | Ainsi, dans cette machine, l'exécution Python/OpenCL est 271 fois plus rapide sur le GPU qu'en Python natif. Si nous désirons "rester" sur le CPU, les résultats sont aussi éloquents. | ||
+ | |||
+ | La version CUDA appelée par <code>./MySteps_6.py -d 0 -g CUDA -s $((2**25)) -c 10 -t 32</code> donne :<code> | ||
+ | Device Selection : 0 | ||
+ | GpuStyle used : CUDA | ||
+ | Size of complex vector : 33554432 | ||
+ | Number of silly calls : 10 | ||
+ | Number of Threads : 32 | ||
+ | Serial compute : 1 | ||
+ | Device #0 of type GPU : NVIDIA GeForce GTX TITAN | ||
+ | Device #1 of type GPU : Quadro K420 | ||
+ | NativeRate: 130602 | ||
+ | Definition of kernel : 1.646 | ||
+ | Synthesis of kernel : 0.000 | ||
+ | Allocation on Host for results : 0.041 | ||
+ | Execution of kernel : 0.869 | ||
+ | Execution of kernel : 0.869 | ||
+ | CUDARate: 12511868 | ||
+ | [ 0.0000000e+00 3.5762787e-07 -5.9604645e-07 ... 2.9802322e-06 | ||
+ | 5.9604645e-07 7.1525574e-07] | ||
+ | 0.016533133 | ||
+ | Results between Native & CUDA seem to be too different! | ||
+ | CUDAvsNative ratio: 95.801504</code> | ||
+ | |||
+ | L'exécution de de <code>./MySteps_6.py -d 4 -g OpenCL -s $((2**25)) -c 10</code> sur sa CPU assez ancienne ne disposant que de 6 coeurs donne : <code> | ||
+ | Device Selection : 4 | ||
+ | GpuStyle used : OpenCL | ||
+ | Size of complex vector : 33554432 | ||
+ | Number of silly calls : 10 | ||
+ | Number of Threads : 1 | ||
+ | Serial compute : 1 | ||
+ | Device #0 from NVIDIA Corporation of type xPU : NVIDIA GeForce GTX TITAN | ||
+ | Device #1 from NVIDIA Corporation of type xPU : Quadro K420 | ||
+ | Device #2 from The pocl project of type xPU : pthread-Intel(R) Xeon(R) CPU E5-2620 0 @ 2.00GHz | ||
+ | Device #3 from Advanced Micro Devices, Inc. of type xPU : Intel(R) Xeon(R) CPU E5-2620 0 @ 2.00GHz | ||
+ | Device #4 from Intel(R) Corporation of type xPU : Intel(R) Xeon(R) CPU E5-2620 0 @ 2.00GHz | ||
+ | NativeRate: 139769 | ||
+ | CPU/GPU selected: Intel(R) Xeon(R) CPU E5-2620 0 @ 2.00GHz | ||
+ | Copy from Host 2 Device : 0.130 | ||
+ | /usr/lib/python3/dist-packages/pyopencl/__init__.py:266: CompilerWarning: Non-empty compiler output encountered. Set the environment variable PYOPENCL_COMPILER_OUTPUT=1 to see more. | ||
+ | warn("Non-empty compiler output encountered. Set the " | ||
+ | Building kernels : 0.199 | ||
+ | Allocation on Device for results : 0.000 | ||
+ | Synthesis of kernel : 0.049 | ||
+ | Execution of kernel : 9.398 | ||
+ | Allocation on Host for results: 0.000 | ||
+ | Copy from Device 2 Host : 0.061 | ||
+ | OpenCLRate: 3280123 | ||
+ | [ 3.8743019e-07 5.9604645e-08 -1.1920929e-07 ... 1.7881393e-07 | ||
+ | 0.0000000e+00 -8.9406967e-07] | ||
+ | 0.038126085 | ||
+ | Results between Native & OpenCL seem to be too different! | ||
+ | OpenCLvsNative ratio: 23.468172 | ||
+ | </code> | ||
+ | |||
+ | Pour une exécution en OpenMP et OpenACC avec les mêmes arguments, nous avons : | ||
+ | * en OpenMP : <code> | ||
+ | 33554432 10 | ||
+ | Norm: 0.00000000e+00 | ||
+ | Elapsed Time: 418.251 | ||
+ | OMP Elapsed Time: 47.402 | ||
+ | NativeRate: 80225 | ||
+ | OMPRate: 707876 | ||
+ | AccRatio: 8.824</code> | ||
+ | * en OpenACC : <code> | ||
+ | 33554432 10 | ||
+ | Norm: 0.00000000e+00 | ||
+ | Elapsed Time: 416.558 | ||
+ | OpenACC Elapsed Time: 12.511 | ||
+ | NativeRate: 80551 | ||
+ | OpenACCRate: 2681906 | ||
+ | ACCRatio: 33.294 | ||
+ | </code> | ||
+ | |||
+ | Ainsi, pour une charge calculatoire de 321 fonctions arithmétiques sur chaque élément du tableau (10x la fonction de Mylq sur chaque élément et l'addition finale) : | ||
+ | * Python/Numpy est presque 2x plus rapide que la C sérielle | ||
+ | * Python/OpenCL/GPU est 8x plus rapide que la C/OpenACC, 271x plus rapide que la Python Numpy | ||
+ | * Python/CUDA/GPU est 3x plus rapide que la C/OpenACC, 10x plus rapide que la Python Numpy | ||
+ | * Python/OpenCL/CPU est presque 3x plus rapide que la C/OpenMP, 23x plus rapide que le Python Numpy | ||
+ | |||
+ | Si nous augmentons la charge individuelle de chaque addition, nous augmentons de manière très significative l'accéleration de calcul via OpenCL. Pour la configuration matérielle ci-dessus, l'accélération frise avec les 1000 pour le GPU et les 23 pour le CPU, en sollicitant 1000 fois la fonction de Mylq (soit 32001 opérations pour chaque somme). | ||
+ | |||
+ | Le tableau suivant montre, pour les 7 implémentations, le gain face à l'implémentation native Python : | ||
+ | ^ Silly Calls ^ C/Serial ^ C/OpenMP ^ C/OpenACC ^ PyCL CPU ^ PyCL GPU ^ PyCUDA 32T ^ | ||
+ | ^ 0| 3.89| 9.39| 1.45| 0.82| 0.60| 0.57| | ||
+ | ^ 1| 0.40| 4.82| 12.95| 15.40| 56.43| 46.08| | ||
+ | ^ 10| 0.46| 4.69| 15.96| 20.30| 393.72| 190.89| | ||
+ | ^ 100| 0.54| 5.18| 17.97| 23.24| 904.59| 294.82| | ||
+ | ^ 1000| 0.52| 5.10| 17.30| 22.40| 967.09| 290.60| | ||
+ | |||
+ | Pour obtenir de tels gains, nous avons la conjonction des 2 facteurs : l'exploitation massive de tous les //cudacores// (les ALU de la GPU) et l'utilisation de toute la bande passante mémoire de la GPU (10x supérieure généralement à la bande passante mémoire). | ||
+ | |||
+ | L'objectif de cette partie TP est donc de retrouver, par soi-même, les facteurs d'accélération de OpenCL pour des charges croissantes par l'application de 0, 1, 10, 100, 1000 fois la fonction de Mylq. | ||
+ | |||
+ | La difficulté viendra du temps d'exécution de la version séquentielle qui atteint son pallier de performances pour une taille croissante de vecteurs assez rapidement et du choix judicieux de l'inhibition de l'exécution séquentielle pour permettre d'atteindre des tailles de vecteurs significatives. | ||
+ | |||
+ | <note warning>**Exercice #4.1 : exécution des différentes implémentations** | ||
+ | - Compilez ''MySteps_6.c'' en ''MySteps_6'' | ||
+ | - Compilez ''MySteps_6_openmp.c'' en ''MySteps_6_openmp'' | ||
+ | - Compilez ''MySteps_6_openacc.c'' en ''MySteps_6_openacc'' | ||
+ | - Compilez ''MySteps_6_openmp.c'' en ''MySteps_6_openmp_NoSerial'', sans exécution séquentielle | ||
+ | - Compilez ''MySteps_6_openacc.c'' en ''MySteps_6_openacc_NoSerial'', sans exécution séquentielle | ||
+ | - Exécutez ''MySteps_6'' sur des tailles de 32768 et 65536, pour les différentes charges ci-dessus | ||
+ | - Relevez les **NativeRate** pour les différentes charges : que constatez vous ? | ||
+ | - Exécutez ''MySteps_6_openmp'' sur des tailles de 32768 et 65536, pour les mêmes charges | ||
+ | - Relevez les **NativeRate**, **OMPRate** pour les différentes charges | ||
+ | - Exécutez ''MySteps_6_openacc'' sur des tailles de 32768 et 65536, pour les mêmes charges | ||
+ | - Relevez les **NativeRate**, **ACCRate** pour les différentes charges | ||
+ | - Exécutez ''MySteps_6.py'' en OpenCL/CPU avec l'Intel sur des tailles de 32768 et 65536 | ||
+ | - Relevez les **NativeRate**, **OpenCLRate** pour les différentes charges | ||
+ | - Exécutez ''MySteps_6.py'' en OpenCL/GPU sur le plus gros GPU sur des tailles de 32768 et 65536 | ||
+ | - Relevez les **NativeRate**, **OpenCLRate** pour les différentes charges | ||
+ | - Exécutez ''MySteps_6.py'' en CUDA/GPU sur le plus gros GPU sur des tailles de 32768 et 65536 | ||
+ | - Relevez les **NativeRate**, **CUDARate** pour les différentes charges | ||
+ | - Placez dans un tableau les différentes valeurs : que constatez-vous ? | ||
+ | - Exécutez ''MySteps_6_openmp_NoSerial'' sur des tailles de 1048576 et 2097152, pour les mêmes charges | ||
+ | - Relevez **OMPRate** pour les différentes charges | ||
+ | - Exécutez ''MySteps_6_openacc_NoSerial'' sur des tailles de 1048576 et 2097152, pour les mêmes charges | ||
+ | - Relevez **ACCRate** pour les différentes charges | ||
+ | - Exécutez ''MySteps_6.py'' en OpenCL Intel sur des tailles de 1048576 et 2097152, avec l'option ''-n'' | ||
+ | - Relevez **OpenCLRate** pour les différentes charges | ||
+ | - Exécutez ''MySteps_6.py'' en OpenCL sur le même GPU sur des tailles de 1048576 et 2097152 | ||
+ | - Relevez **OpenCLRate** pour les différentes charges | ||
+ | - Exécutez ''MySteps_6.py'' en CUDA/GPU sur le même GPU sur des tailles de 1048576 et 2097152 | ||
+ | - Relevez **CUDARate** pour les différentes charges | ||
+ | - Placez dans un tableau les différentes valeurs : que constatez-vous comme gain en performance en OpenCL ? | ||
+ | </note> | ||
+ | |||
+ | Vous pouvez juger, du fait de lancement pour des tailles doublées, et pour les différentes charges, que vous avez obtenu des optimums pour certaines implémentations, mais pas encore pour d'autres, notamment OpenCL et CUDA. | ||
+ | |||
+ | Si la performance pour une taille de 2097152 est moins de 5% supérieure à la performance pour une taille de 1048576, vous pouvez considérer que vous avez atteint le quasi-optimum de performance. L'objectif est d'atteindre cette limite. | ||
+ | |||
+ | <note warning>**Exercice #4.2 : exploration de la meilleure performance OpenCL et CUDA** | ||
+ | - Exécutez ''MySteps_6.py'' en OpenCL Intel sur des tailles croissantes avec l'option ''-n'' | ||
+ | - Relevez **OpenCLRate** pour les différentes charges | ||
+ | - Exécutez ''MySteps_6.py'' en OpenCL sur le même GPU sur des tailles de 1048576 et 2097152 | ||
+ | - Relevez **OpenCLRate** pour les différentes charges | ||
+ | - Exécutez ''MySteps_6.py'' en CUDA/GPU sur le même GPU sur des tailles de 1048576 et 2097152 | ||
+ | - Relevez **CUDARate** pour les différentes charges | ||
+ | - Placez dans un tableau les différentes valeurs : que constatez-vous comme gain en performance en OpenCL ? | ||
+ | </note> | ||
+ | |||
===== Implémenter une fonction "coûteuse", la Transformée de Fourier ===== | ===== Implémenter une fonction "coûteuse", la Transformée de Fourier ===== | ||
Ligne 891: | Ligne 1091: | ||
Il sera alors possible d'estimer l'erreur numérique à ce calcul. | Il sera alors possible d'estimer l'erreur numérique à ce calcul. | ||
- | <note warning>**Exercice #4.1 : implémentation Python "naïve"** | + | <note warning>**Exercice #5.1 : implémentation Python "naïve"** |
- Modifiez ''MyDFT_1.py'' suivant les 7 spécifications ci-dessus | - Modifiez ''MyDFT_1.py'' suivant les 7 spécifications ci-dessus | ||
- Exécutez le programme pour une taille de **16** et contrôler la cohérence | - Exécutez le programme pour une taille de **16** et contrôler la cohérence | ||
Ligne 924: | Ligne 1124: | ||
- comparer les résultats entre les deux avec ''linalg.norm'' | - comparer les résultats entre les deux avec ''linalg.norm'' | ||
- | <note warning>**Exercice #4.2 : implémentation Python Numpy** | + | <note warning>**Exercice #5.2 : implémentation Python Numpy** |
- Copiez le programme ''MyDFT_1.py'' en ''MyDFT_2.py'' | - Copiez le programme ''MyDFT_1.py'' en ''MyDFT_2.py'' | ||
- Modifiez ''MyDFT_2.py'' suivant les 7 spécifications ci-dessus | - Modifiez ''MyDFT_2.py'' suivant les 7 spécifications ci-dessus | ||
Ligne 951: | Ligne 1151: | ||
- changer le domaine d'itération pour la boucle : ''range()'' par ''numba.prange()'' | - changer le domaine d'itération pour la boucle : ''range()'' par ''numba.prange()'' | ||
- | <note warning>**Exercice #4.3 : implémentation Python Numpy** | + | <note warning>**Exercice #5.3 : implémentation Python Numpy** |
- Copiez le programme ''MyDFT_2.py'' en ''MyDFT_3.py'' et exploitez ce dernier | - Copiez le programme ''MyDFT_2.py'' en ''MyDFT_3.py'' et exploitez ce dernier | ||
- Copiez la fonction ''NumpyDFT'' en ''NumbaDFT'' | - Copiez la fonction ''NumpyDFT'' en ''NumbaDFT'' | ||
Ligne 981: | Ligne 1181: | ||
Pour l'implémentation OpenCL, la version "naïve" de l'implémentation va servir. Pour cela, il suffit de reprendre la définition de la méthode naïve et de l'implémenter en C dans un noyau OpenCL. A noter que Pi n'étant dans une variable définie, il faut explicitement la détailler dans le noyau OpenCL. Autre détail important : le //cast//. De manière a éviter tout effet de bord, il est fortement recommandé de //caster// les opérations dans la précision flottante souhaitée pour des opérations sur des indices entiers. | Pour l'implémentation OpenCL, la version "naïve" de l'implémentation va servir. Pour cela, il suffit de reprendre la définition de la méthode naïve et de l'implémenter en C dans un noyau OpenCL. A noter que Pi n'étant dans une variable définie, il faut explicitement la détailler dans le noyau OpenCL. Autre détail important : le //cast//. De manière a éviter tout effet de bord, il est fortement recommandé de //caster// les opérations dans la précision flottante souhaitée pour des opérations sur des indices entiers. | ||
- | <note warning>**Exercice #4.4 : implémentation Python OpenCL** | + | <note warning>**Exercice #5.4 : implémentation Python OpenCL** |
- Copiez le programme ''MyDFT_3.py'' en ''MyDFT_4.py'' et exploitez ce dernier | - Copiez le programme ''MyDFT_3.py'' en ''MyDFT_4.py'' et exploitez ce dernier | ||
- Copiez la fonction python ''OpenCLAddition'' en ''OpenCLDFT'' | - Copiez la fonction python ''OpenCLAddition'' en ''OpenCLDFT'' | ||
Ligne 1056: | Ligne 1256: | ||
* modifier les vecteurs en sortie (2 vecteurs) | * modifier les vecteurs en sortie (2 vecteurs) | ||
- | <note warning>**Exercice #4.5 : implémentation Python CUDA** | + | <note warning>**Exercice #5.5 : implémentation Python CUDA** |
- Copiez le programme ''MyDFT_4.py'' en ''MyDFT_5.py'' et exploitez ce dernier | - Copiez le programme ''MyDFT_4.py'' en ''MyDFT_5.py'' et exploitez ce dernier | ||
- Copiez la fonction python ''CUDAAddition'' en ''CUDADFT'' | - Copiez la fonction python ''CUDAAddition'' en ''CUDADFT'' | ||
Ligne 1165: | Ligne 1365: | ||
L'objectif est donc de reprendre notre exemple le plus abouti de notre DFT et d'y ajouter ces éléments. Pour cela, les programmes ''PiXPU.py'' et ''TrouNoir.py'' vont être explorés pour voir comment faire. | L'objectif est donc de reprendre notre exemple le plus abouti de notre DFT et d'y ajouter ces éléments. Pour cela, les programmes ''PiXPU.py'' et ''TrouNoir.py'' vont être explorés pour voir comment faire. | ||
- | <note warning>**Exercice #5.1 : exploration de PiXPU.py** | + | <note warning>**Exercice #6.1 : exploration de PiXPU.py** |
- Identifiez les lignes correspondant aux paramétrages par défaut | - Identifiez les lignes correspondant aux paramétrages par défaut | ||
- Identifiez les lignes sur la découverte des périphériques OpenCL | - Identifiez les lignes sur la découverte des périphériques OpenCL | ||
Ligne 1183: | Ligne 1383: | ||
- Sélectionner une exécution sous OpenCL ou CUDA avec l'option ''-g'' | - Sélectionner une exécution sous OpenCL ou CUDA avec l'option ''-g'' | ||
- | <note warning>**Exercice #5.2 : modification du programme ''MyDFT_6.py''** | + | <note warning>**Exercice #6.2 : modification du programme ''MyDFT_6.py''** |
- Supprimer la sélection initiale d'argument | - Supprimer la sélection initiale d'argument | ||
- Inhiber pour l'instant l'exécution des fonctions | - Inhiber pour l'instant l'exécution des fonctions | ||
Ligne 1224: | Ligne 1424: | ||
</code> | </code> | ||
- | <note warning>**Exercice #5.3 : modification du programme ''MyDFT_7.py''** | + | <note warning>**Exercice #6.3 : modification du programme ''MyDFT_7.py''** |
- Libérez pour l'appel à la fonction ''OpenCLDFT'' | - Libérez pour l'appel à la fonction ''OpenCLDFT'' | ||
- Rajoutez le test exploitant la sélection OpenCL ou CUDA | - Rajoutez le test exploitant la sélection OpenCL ou CUDA | ||
Ligne 1402: | Ligne 1602: | ||
<note warning> | <note warning> | ||
- | **Exercice #6.1 : éditez le source du programme ''xGEMM.c'' et repérez les éléments suivants** | + | **Exercice #7.1 : éditez le source du programme ''xGEMM.c'' et repérez les éléments suivants** |
* Identifiez dans ''Makefile'' quelles directives (précédées par ''-D'') sont associées aux différentes implémentations | * Identifiez dans ''Makefile'' quelles directives (précédées par ''-D'') sont associées aux différentes implémentations | ||
Ligne 1671: | Ligne 1871: | ||
<note warning> | <note warning> | ||
- | **Exercice #6.2 : lancez les ''xGEMM_<precision>_<implementation>'' avec une taille de 1000** | + | **Exercice #7.2 : lancez les ''xGEMM_<precision>_<implementation>'' avec une taille de 1000** |
* Variez le nombre d'itérations pour obtenir une durée d'exécution d'une dizaine de secondes ? | * Variez le nombre d'itérations pour obtenir une durée d'exécution d'une dizaine de secondes ? | ||
Ligne 1690: | Ligne 1890: | ||
<note warning> | <note warning> | ||
- | **Exercice #6.3 : lancez les programmes précédents pour différentes tailles** | + | **Exercice #7.3 : lancez les programmes précédents pour différentes tailles** |
* Diminuez la taille aux valeurs suivantes ''125'', ''250'', ''500'' et exécutez les programmes | * Diminuez la taille aux valeurs suivantes ''125'', ''250'', ''500'' et exécutez les programmes | ||
Ligne 1725: | Ligne 1925: | ||
L'objectif est de "jouer" le [[https://www.tensorflow.org/tutorials/images/cnn|tutoriel]] exploitant la base d'images CIFAR10 pour un apprentissage convolutif. | L'objectif est de "jouer" le [[https://www.tensorflow.org/tutorials/images/cnn|tutoriel]] exploitant la base d'images CIFAR10 pour un apprentissage convolutif. | ||
- | <note warning>**Exercice #7.1 :** | + | <note warning>**Exercice #8.1 :** |
- chargez l'environnement conda | - chargez l'environnement conda | ||
- préparez la variable d'environnement ''TIME'' | - préparez la variable d'environnement ''TIME'' | ||
Ligne 1748: | Ligne 1948: | ||
En regardant l'activité du GPU, il apparaît que le gain est substanciel par rapport à une "petite" configuration GPU. Cependant, la nature du réseau créé n'exploitait pas de manière optimale la GPU par rapport à la CPU. Une petite modification de notre réseau va permettre de mettre cela en évidence, en modifiant le nombre de poids d'une des couches neuronales. | En regardant l'activité du GPU, il apparaît que le gain est substanciel par rapport à une "petite" configuration GPU. Cependant, la nature du réseau créé n'exploitait pas de manière optimale la GPU par rapport à la CPU. Une petite modification de notre réseau va permettre de mettre cela en évidence, en modifiant le nombre de poids d'une des couches neuronales. | ||
- | <note warning>**Exercice #7.2 :** | + | <note warning>**Exercice #8.2 :** |
- changez **64** en **65536** dans ''model.add(layers.Dense(64, activation='relu'))'' | - changez **64** en **65536** dans ''model.add(layers.Dense(64, activation='relu'))'' | ||
- supprimez la référence à ''CUDA_VISIBLE_DEVICES'' avec ''export -n CUDA_VISIBLE_DEVICES'' | - supprimez la référence à ''CUDA_VISIBLE_DEVICES'' avec ''export -n CUDA_VISIBLE_DEVICES'' | ||
Ligne 1771: | Ligne 1971: | ||
Le code source est accessible à l'adresse : https://www.r-ccs.riken.jp/labs/cbrt/download/genesis-version-1-5/ | Le code source est accessible à l'adresse : https://www.r-ccs.riken.jp/labs/cbrt/download/genesis-version-1-5/ | ||
- | <note warning>**Exercice #8.1 : Récupérez et compilez le code suivant la documentation fournie** | + | <note warning>**Exercice #9.1 : Récupérez et compilez le code suivant la documentation fournie** |
* Lisez la [[https://www.r-ccs.riken.jp/labs/cbrt/installation/|documentation]] d'installation | * Lisez la [[https://www.r-ccs.riken.jp/labs/cbrt/installation/|documentation]] d'installation | ||
* Placez les sources dans le dossier ''/local/$USER/GENESIS'' créé pour l'occasion | * Placez les sources dans le dossier ''/local/$USER/GENESIS'' créé pour l'occasion | ||
Ligne 1790: | Ligne 1990: | ||
Pour finir, dans comme ce programme est "aussi" //gépufié// (porté sur GPU), il risque d'y avoir un goulet d'étranglement pour l'accès au GPU pour les 64 tâches simultanées. Ainsi, les programmes "fortement" parallélisés exigent de choisir judicieusement les différents paramètres de parallélisation tout comme nous avons vue que, pour les GPU, il fallait découper la tâche en un nombre optimal de sous-tâches. | Pour finir, dans comme ce programme est "aussi" //gépufié// (porté sur GPU), il risque d'y avoir un goulet d'étranglement pour l'accès au GPU pour les 64 tâches simultanées. Ainsi, les programmes "fortement" parallélisés exigent de choisir judicieusement les différents paramètres de parallélisation tout comme nous avons vue que, pour les GPU, il fallait découper la tâche en un nombre optimal de sous-tâches. | ||
- | <note warning>**Exercice #8.2 : Exécutez l'exemple ''alad_water''** | + | <note warning>**Exercice #9.2 : Exécutez l'exemple ''alad_water''** |
* Récupérez [[http://www.cbp.ens-lyon.fr/emmanuel.quemener/documents/alad_water.tgz|l'exemple d'exécution]] | * Récupérez [[http://www.cbp.ens-lyon.fr/emmanuel.quemener/documents/alad_water.tgz|l'exemple d'exécution]] | ||
* Décompressez l'archive dans ''/local/$USER/GENESIS'' | * Décompressez l'archive dans ''/local/$USER/GENESIS'' | ||
Ligne 1818: | Ligne 2018: | ||
Nous allons tenter de reproduire une [[https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/gromacs/|expérience de Nvidia]] vantant l'efficacité des GPGPU pour le logiciel de [[https://fr.wikipedia.org/wiki/Dynamique_mol%C3%A9culaire|dynamique moléculaire]] [[http://www.gromacs.org/|Gromacs]]. | Nous allons tenter de reproduire une [[https://www.nvidia.com/en-us/data-center/gpu-accelerated-applications/gromacs/|expérience de Nvidia]] vantant l'efficacité des GPGPU pour le logiciel de [[https://fr.wikipedia.org/wiki/Dynamique_mol%C3%A9culaire|dynamique moléculaire]] [[http://www.gromacs.org/|Gromacs]]. | ||
- | <note warning>**Exercice #9.1 : appliquez la "recette" de Nvidia** | + | <note warning>**Exercice #10.1 : appliquez la "recette" de Nvidia** |
* La documentation offre ceci : | * La documentation offre ceci : | ||
- récupérez le source | - récupérez le source | ||
Ligne 1833: | Ligne 2033: | ||
En cas de difficultés, appliquez la [[formation:insa2020gpu:insa2020gromacs4buster|recette de Gromacs pour Debian Buster]] ;-) | En cas de difficultés, appliquez la [[formation:insa2020gpu:insa2020gromacs4buster|recette de Gromacs pour Debian Buster]] ;-) | ||
- | <note warning>**Exercice #9.2 : Exécutez l'exemple ''1536''** | + | <note warning>**Exercice #10.2 : Exécutez l'exemple ''1536''** |
* Quel ''Elapsed Time'' avez-vous pour l'exécution sur GPU (et CPU) ? | * Quel ''Elapsed Time'' avez-vous pour l'exécution sur GPU (et CPU) ? | ||
* Quel ''Elapsed Time'' avez-vous pour l'exécution uniquement sur CPU ? | * Quel ''Elapsed Time'' avez-vous pour l'exécution uniquement sur CPU ? |