Différences
Ci-dessous, les différences entre deux révisions de la page.
| Les deux révisions précédentes Révision précédente | |||
|
formation:insa2024gpu [2024/11/16 12:14] equemene [Première exploration de l'association Python et OpenCL] |
formation:insa2024gpu [2024/11/16 12:16] (Version actuelle) equemene |
||
|---|---|---|---|
| Ligne 1361: | Ligne 1361: | ||
| Les méthodes présentées dans la suite ont l'avantage d'être systématiques et donc pourront être exploitées dans n'importe quel programme PyOpenCL pour PyCUDA. | Les méthodes présentées dans la suite ont l'avantage d'être systématiques et donc pourront être exploitées dans n'importe quel programme PyOpenCL pour PyCUDA. | ||
| - | A cela s'ajoute également la possibilité d'ajouter des options au lancement du programme, notamment pour spécifier sa GPU. Dans les programmes de l'archive **bench4gpu** en Python, la majorité dispose d'une option ''-h'' permettant de voir les options ainsi que de lister les périphériques OpenCL ou CUDA détectés. | + | A cela s'ajoute également la possibilité d'ajouter des options au lancement du programme, notamment pour spécifier sa GPU. Dans les programmes de l'archive **bench4xpu** en Python, la majorité dispose d'une option ''-h'' permettant de voir les options ainsi que de lister les périphériques OpenCL ou CUDA détectés. |
| 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. | ||
| Ligne 1531: | Ligne 1531: | ||
| GPU selected Quadro K420 | GPU selected Quadro K420 | ||
| Traceback (most recent call last): | Traceback (most recent call last): | ||
| - | File "/home/equemene/bench4gpu/ETSN/./MyDFT_8.py", line 356, in <module> | + | File "/home/equemene/bench4xpu/ETSN/./MyDFT_8.py", line 356, in <module> |
| k_np,l_np=CUDADFT(a_np,b_np,Device) | k_np,l_np=CUDADFT(a_np,b_np,Device) | ||
| - | File "/home/equemene/bench4gpu/ETSN/./MyDFT_8.py", line 162, in CUDADFT | + | File "/home/equemene/bench4xpu/ETSN/./MyDFT_8.py", line 162, in CUDADFT |
| mod = SourceModule(""" | mod = SourceModule(""" | ||
| File "/usr/lib/python3/dist-packages/pycuda/compiler.py", line 290, in __init__ | File "/usr/lib/python3/dist-packages/pycuda/compiler.py", line 290, in __init__ | ||
| Ligne 1595: | Ligne 1595: | ||
| ==== Exploitations de xGEMM ==== | ==== Exploitations de xGEMM ==== | ||
| - | Le dossier ''bench4gpu/BLAS/xGEMM'' contient peu de fichiers dont les importants sont : un unique programme source, ''xGEMM.c'', et un fichier de construction, ''Makefile''. C'est ce fichier qui va ''construire'' tous les exécutables d'un coup, à la fois pour les différentes implémentations de BLAS, mais aussi pour les deux précisions **SP** (simple précision sur 32 bits) et **DP** (double précision sur 64 bits). | + | Le dossier ''bench4xpu/BLAS/xGEMM'' contient peu de fichiers dont les importants sont : un unique programme source, ''xGEMM.c'', et un fichier de construction, ''Makefile''. C'est ce fichier qui va ''construire'' tous les exécutables d'un coup, à la fois pour les différentes implémentations de BLAS, mais aussi pour les deux précisions **SP** (simple précision sur 32 bits) et **DP** (double précision sur 64 bits). |
| === Le source === | === Le source === | ||
| Ligne 1695: | Ligne 1695: | ||
| Error 0.0000000000 | Error 0.0000000000 | ||
| - | root@opencluster2:/local/tests/bench4gpu/BLAS/xGEMM# ./xGEMM_SP_thunking 1000 1000 | + | root@opencluster2:/local/tests/bench4xpu/BLAS/xGEMM# ./xGEMM_SP_thunking 1000 1000 |
| Using CuBLAS/Thunking: 1000 iterations for 1000x1000 matrix | Using CuBLAS/Thunking: 1000 iterations for 1000x1000 matrix | ||
| Ligne 1709: | Ligne 1709: | ||
| Dans le premier terminal qui nous sert à l'exécution des programmes, nous avons : | Dans le premier terminal qui nous sert à l'exécution des programmes, nous avons : | ||
| <code> | <code> | ||
| - | root@opencluster2:/local/tests/bench4gpu/BLAS/xGEMM# ./xGEMM_SP_cublas 1000 100000 | + | root@opencluster2:/local/tests/bench4xpu/BLAS/xGEMM# ./xGEMM_SP_cublas 1000 100000 |
| Using CuBLAS: 100000 iterations for 1000x1000 matrix | Using CuBLAS: 100000 iterations for 1000x1000 matrix | ||
| Ligne 1820: | Ligne 1820: | ||
| 1 warning generated. | 1 warning generated. | ||
| ./xGEMM_SP_clblas: symbol lookup error: /root/.cache/pocl/kcache/DO/MOINFKPIFEHKLMFJBAMOHIHHPGDFMKFNMCFAG/sgemm_Col_NN_B0_MX032_NX032_KX08/16-16-1/sgemm_Col_NN_B0_MX032_NX032_KX08.so: undefined symbol: mem_fence | ./xGEMM_SP_clblas: symbol lookup error: /root/.cache/pocl/kcache/DO/MOINFKPIFEHKLMFJBAMOHIHHPGDFMKFNMCFAG/sgemm_Col_NN_B0_MX032_NX032_KX08/16-16-1/sgemm_Col_NN_B0_MX032_NX032_KX08.so: undefined symbol: mem_fence | ||
| - | root@opencluster2:/local/tests/bench4gpu/BLAS/xGEMM# ./xGEMM_SP_clblas 1000 10 1 0 | + | root@opencluster2:/local/tests/bench4xpu/BLAS/xGEMM# ./xGEMM_SP_clblas 1000 10 1 0 |
| Using CLBLAS: 10 iterations for 1000x1000 matrix on (1,0) | Using CLBLAS: 10 iterations for 1000x1000 matrix on (1,0) | ||
| Device (1,0): pthread-Intel(R) Xeon(R) CPU E5-2637 v4 @ 3.50GHz | Device (1,0): pthread-Intel(R) Xeon(R) CPU E5-2637 v4 @ 3.50GHz | ||
| Ligne 2080: | Ligne 2080: | ||
| ==== Les programmes utilisés : OpenCL, CUDA ==== | ==== Les programmes utilisés : OpenCL, CUDA ==== | ||
| - | Le dossier ''bench4gpu'' contient de nombreuses implémentations de ce calcul élémentaire. Nous nous focaliserons sur 2 d'entre eux : | + | Le dossier ''bench4xpu'' contient de nombreuses implémentations de ce calcul élémentaire. Nous nous focaliserons sur 2 d'entre eux : |
| - | * ''bench4gpu/Pi/OpenCL/PiOpenCL.c'' | + | * ''bench4xpu/Pi/OpenCL/PiOpenCL.c'' |
| - | * ''bench4gpu/Pi/XPU/PiXPU.py'' | + | * ''bench4xpu/Pi/XPU/PiXPU.py'' |
| ==== Implémentation C/OpenCL ==== | ==== Implémentation C/OpenCL ==== | ||
| - | Le programme ''bench4gpu/Pi/C/OpenCL/PiOpenCL.c'' est une implémentation en C "pur" : il permet de juger de la difficulté d'appropriation de OpenCL. En effet, une grande partie du code est destinée à définir quelle plateforme et quel périphérique utiliser, placer les données à traîter. | + | Le programme ''bench4xpu/Pi/C/OpenCL/PiOpenCL.c'' est une implémentation en C "pur" : il permet de juger de la difficulté d'appropriation de OpenCL. En effet, une grande partie du code est destinée à définir quelle plateforme et quel périphérique utiliser, placer les données à traîter. |
| <note warning> | <note warning> | ||
| Ligne 2282: | Ligne 2282: | ||
| Dans l'exemple précédent, nous avons exploité un programme en C "pur". Nous avons vu comment s'intégrait la portion de code (le noyau ou le //kernel//) qui était exécutée, soit par les GPU, soit par différentes implémentations pour CPU de OpenCL. Nous avons aussi pu "juger" du côté assez "compliqué" de découvertes des périphériques et des mécanismes d'entrée/sortie. | Dans l'exemple précédent, nous avons exploité un programme en C "pur". Nous avons vu comment s'intégrait la portion de code (le noyau ou le //kernel//) qui était exécutée, soit par les GPU, soit par différentes implémentations pour CPU de OpenCL. Nous avons aussi pu "juger" du côté assez "compliqué" de découvertes des périphériques et des mécanismes d'entrée/sortie. | ||
| - | Avec le programme ''PiXPU.py'', situé dans le dossier ''bench4gpu/Pi/XPU'' nous allons pouvoir exploiter les périphériques de calcul plus simplement, et bénéficier de toutes les facilités d'un langage évolué. | + | Avec le programme ''PiXPU.py'', situé dans le dossier ''bench4xpu/Pi/XPU'' nous allons pouvoir exploiter les périphériques de calcul plus simplement, et bénéficier de toutes les facilités d'un langage évolué. |
| Regardons d'abord sa sortie lorsqu'il est sollicité avec l'option ''-h'' :<code> | Regardons d'abord sa sortie lorsqu'il est sollicité avec l'option ''-h'' :<code> | ||
| - | $ cd $HOME/bench4gpu/Pi/XPU | + | $ cd $HOME/bench4xpu/Pi/XPU |
| $ python3 PiXPU.py -h | $ python3 PiXPU.py -h | ||
| PiXPU.py -o (Out of Core Metrology) -c (Print Curves) -k (Case On IfThen) -d <DeviceId> -g <CUDA/OpenCL> -i <Iterations> -b <BlocksBegin> -e <BlocksEnd> -s <BlocksStep> -f <ThreadsFirst> -l <ThreadsLast> -t <ThreadssTep> -r <RedoToImproveStats> -m <SHR3/CONG/MWC/KISS> -v <INT32/INT64/FP32/FP64> | PiXPU.py -o (Out of Core Metrology) -c (Print Curves) -k (Case On IfThen) -d <DeviceId> -g <CUDA/OpenCL> -i <Iterations> -b <BlocksBegin> -e <BlocksEnd> -s <BlocksStep> -f <ThreadsFirst> -l <ThreadsLast> -t <ThreadssTep> -r <RedoToImproveStats> -m <SHR3/CONG/MWC/KISS> -v <INT32/INT64/FP32/FP64> | ||
| Ligne 2459: | Ligne 2459: | ||
| Le code de test **Pi Monte Carlo** avait pour avantage de ne solliciter que très peu la mémoire. La seule mémoire sollicitée sur les périphériques se résumait au nombre de tirages dans le quadrant de cercle. De plus, l'opération de parallélisation divisait un nombre total d'itérations en un nombre équivalent pour chaque //work item//, chaque //block// ou chaque //thread//. Chaque calcul était indépendant des autres : nous avions non seulement un **code ALU** (ne sollicitant que les unités de traitement) mais aussi un **code gros grain** (indépendance des calculs). | Le code de test **Pi Monte Carlo** avait pour avantage de ne solliciter que très peu la mémoire. La seule mémoire sollicitée sur les périphériques se résumait au nombre de tirages dans le quadrant de cercle. De plus, l'opération de parallélisation divisait un nombre total d'itérations en un nombre équivalent pour chaque //work item//, chaque //block// ou chaque //thread//. Chaque calcul était indépendant des autres : nous avions non seulement un **code ALU** (ne sollicitant que les unités de traitement) mais aussi un **code gros grain** (indépendance des calculs). | ||
| - | Dans le code ''NBody.py'', situé dans ''bench4gpu/NBody'', le principe est de déterminer dans un modèle newtonien la position et la vitesse de chaque particule en intéraction avec toutes les autres. A chaque itération, chaque position et chaque vitesse de chaque particule vont être modifiées par la proximité de toutes les autres : nous disposons d'un **code grain fin**. De plus, la détermination de la vitesse et la position de chaque particule à chaque instant nécessite un accès à la position de chacune des autres : un accès permanent à la mémoire est donc indispensable. | + | Dans le code ''NBody.py'', situé dans ''bench4xpu/NBody'', le principe est de déterminer dans un modèle newtonien la position et la vitesse de chaque particule en intéraction avec toutes les autres. A chaque itération, chaque position et chaque vitesse de chaque particule vont être modifiées par la proximité de toutes les autres : nous disposons d'un **code grain fin**. De plus, la détermination de la vitesse et la position de chaque particule à chaque instant nécessite un accès à la position de chacune des autres : un accès permanent à la mémoire est donc indispensable. |
| Le programme ''NBody.py'' n'est disponible qu'en OpenCL. Son invocation avec l'option ''-h'' offre la sortie suivante : | Le programme ''NBody.py'' n'est disponible qu'en OpenCL. Son invocation avec l'option ''-h'' offre la sortie suivante : | ||
| - | <code>$ cd $HOME/bench4gpu/NBody | + | <code>$ cd $HOME/bench4xpu/NBody |
| $ python3 NBody.py -h | $ python3 NBody.py -h | ||
| NBody.py -h [Help] -r [InitialRandom] -g [OpenGL] -e [VirielStress] -o [Verbose] -p [Potential] -x <None|NegExp|CorRad> -d <DeviceId> -n <NumberOfParticules> -i <Iterations> -z <SizeOfBoxOrBall> -v <Velocity> -s <Step> -b <Ball|Box> -m <ImplicitEuler|RungeKutta|ExplicitEuler|Heun> -t <FP32|FP64> | NBody.py -h [Help] -r [InitialRandom] -g [OpenGL] -e [VirielStress] -o [Verbose] -p [Potential] -x <None|NegExp|CorRad> -d <DeviceId> -n <NumberOfParticules> -i <Iterations> -z <SizeOfBoxOrBall> -v <Velocity> -s <Step> -b <Ball|Box> -m <ImplicitEuler|RungeKutta|ExplicitEuler|Heun> -t <FP32|FP64> | ||