Alliance Doc

OpenACC Tutorial - Adding directives/fr: Difference between revisions

Page Discussion

OpenACC Tutorial - Adding directives/fr (view source)

Revision as of 19:37, 8 June 2023

1,813 bytes added ,  1 year ago
Updating to match new version of source page
No edit summary
(Updating to match new version of source page)
 
(49 intermediate revisions by 4 users not shown)
Line 13: Line 13:


== Transfert vers un processeur graphique (GPU) ==  
== Transfert vers un processeur graphique (GPU) ==  
Avant de porter du code sur un GPU, il faut savoir qu'ils ne partagent pas la même mémoire que le processeur (CPU);  autrement dit, le GPU n'a pas un accès direct à la mémoire de départ. Cette dernière est en général plus grande, mais plus lente que la mémoire du GPU. Pour pouvoir utiliser un GPU, les données doivent passer par le bus PCI, dont la bande passante est moins grande que celles du CPU et du GPU. Il est donc de la plus haute importance de bien gérer les transferts entre la mémoire de départ et le GPU.  En anglais, ce processus s'appelle ''offloading''.
Avant de porter du code sur un GPU, il faut savoir que ceux-ci ne partagent pas la même mémoire que le CPU de l'hôte.
* la mémoire de l'hôte est en général plus grande, mais plus lente que la mémoire du GPU;
* un GPU n'a pas d'accès direct à la mémoire de l'hôte;
* pour pouvoir utiliser un GPU, les données doivent passer par le bus PCI, dont la bande passante est moins grande que celles du CPU et du GPU;
* il est donc de la plus haute importance de bien gérer les transferts entre la mémoire de départ et le GPU.  En anglais, ce processus s'appelle ''offloading''.


==Directives OpenACC== <!--T:4-->
==Directives OpenACC== <!--T:4-->
Les directives OpenAcc sont semblables aux directives OpenMP. En C/C++, ce sont des <tt>pragmas</tt>  et en Fortran, des commentaires. L'emploi de directives comporte plusieurs avantages. Premièrement, puisque le code est peu affecté, les modifications peuvent se faire de manière incrémentale, un <tt>pragma</tt> à la fois;  ceci est particulièrement utile pour le débogage puisqu'il est ainsi facile d'identifier le changement précis qui crée le bogue. Deuxièmement, OpenACC peut être désactivé au moment de la compilation;  les <tt>pragmas</tt> sont alors vus comme étant des commentaires et ne sont pas considérés par le compilateur, ce qui permet de compiler une version accélérée et une version normale à partir du même code source. Troisièmement, comme le compilateur fait tout le travail de transfert, le même code peut être compilé pour différents types d'accélérateurs, que ce soit un GPU, un Xeon Phi (MIC) ou un CPU;    ainsi, un changement du matériel exigera simplement la mise à jour du compilateur, sans modification au code.  
Les directives OpenAcc sont semblables aux directives [[OpenMP/fr|OpenMP]]. En C/C++, ce sont des énoncés <tt>pragmas</tt>  et en Fortran, des commentaires. L'emploi de directives comporte plusieurs avantages&nbsp;:
* Premièrement, puisque le code est peu affecté, les modifications peuvent se faire de manière incrémentale, un <tt>pragma</tt> à la fois;  ceci est particulièrement utile pour le débogage puisqu'il est ainsi facile d'identifier le changement précis qui crée le bogue.
* Deuxièmement, OpenACC peut être désactivé au moment de la compilation;  les <tt>pragmas</tt> sont alors vus comme étant des commentaires et ne sont pas considérés par le compilateur, ce qui permet de compiler une version accélérée et une version normale à partir du même code source.
* Troisièmement, comme le compilateur fait tout le travail de transfert, le même code peut être compilé pour différents types d'accélérateurs, que ce soit un GPU ou des instructions SIMD sur un CPU;    ainsi, un changement du matériel exigera simplement la mise à jour du compilateur, sans modification au code.  


Le code de notre exemple contient deux boucles&nbsp;: la première initialise deux vecteurs et la seconde effectue une opération de [https://fr.wikipedia.org/wiki/Basic_Linear_Algebra_Subprograms niveau 1] d'addition des vecteurs.  
Le code de notre exemple contient deux boucles&nbsp;: la première initialise deux vecteurs et la seconde effectue une opération de [https://fr.wikipedia.org/wiki/Basic_Linear_Algebra_Subprograms niveau 1] d'addition des vecteurs.  


{| class="wikitable" width="100%"
{| class="wikitable" width="100%"
Line 49: Line 56:
</syntaxhighlight>
</syntaxhighlight>
|}
|}
Dans les deux cas, le compilateur identifie deux noyaux (''kernels''). En C/C++, les deux noyaux sont à l'intérieur de chaque boucle. En Fortran, les noyaux sont à l'intérieur de la première boucle et à l'intérieur de la boucle implicite effectuée lors d'une opération sur des tableaux.
Dans les deux cas, le compilateur identifie deux noyaux (''kernels'')&nbsp;:
* en C/C++, les deux noyaux sont à l'intérieur de chaque boucle.;
* en Fortran, les noyaux sont à l'intérieur de la première boucle et à l'intérieur de la boucle implicite effectuée lors d'une opération sur des tableaux.


Remarquez que le bloc OpenACC est délimité en C/C++ par des accolades; en Fortran, le commentaire est placé une fois au début et une dernière fois à la fin, avec l'ajout cette fois de <tt>end</tt>.
Remarquez que le bloc OpenACC est délimité en C/C++ par des accolades; en Fortran, le commentaire est placé une fois au début et une dernière fois à la fin, avec l'ajout cette fois de <tt>end</tt>.


=== Loop et kernels ===
=== Boucles et noyaux ===
Quand le compilateur lit la directive OpenACC <tt>kernels</tt>, il analyse le code pour identifier les sections pouvant être parallélisées. Ceci correspond souvent au corps d'une boucle. Dans ce cas, le compilateur délimite le début et la fin du corps du code avec la fonction ''kernel''. Les appels à cette fonction ne seront pas affectés par les autres appels. La fonction est compilée et peut ensuite être exécutée sur un accélérateur. Comme chaque appel est indépendant, chacun des milliers de cœurs de l'accélérateur peut exécuter la fonction en parallèle pour un index spécifique.
 
Quand le compilateur lit la directive OpenACC <tt>kernels</tt>, il analyse le code pour identifier les sections pouvant être parallélisées.
Ceci correspond souvent au corps d'une boucle qui a des itérations indépendantes.
Dans ce cas, le compilateur délimite le début et la fin du corps du code avec la fonction [https://en.wikipedia.org/wiki/Compute_kernel ''kernel''].
Les appels à cette fonction ne seront pas affectés par les autres appels.
La fonction est compilée et peut ensuite être exécutée sur un accélérateur.
Comme chaque appel est indépendant, chacun des centaines de cœurs de l'accélérateur peut exécuter la fonction en parallèle pour un index spécifique.
 
{| class="wikitable" width="100%"
{| class="wikitable" width="100%"
|-
|-
! LOOP !! KERNELS
! BOUCLE !! KERNEL
|-
|-
| <syntaxhighlight lang="cpp" line>
| <syntaxhighlight lang="cpp" line>
Line 65: Line 81:
}
}
</syntaxhighlight> || <syntaxhighlight lang="cpp" line>
</syntaxhighlight> || <syntaxhighlight lang="cpp" line>
void loopBody(A,B,C,i)
void kernelName(A, B, C, i)
{
{
   C[i] = A[i] + B[i];
   C[i] = A[i] + B[i];
Line 71: Line 87:
</syntaxhighlight>
</syntaxhighlight>
|-
|-
|Calcule de 0-N, dans l'ordre || Chaque cœur de calcul traite une valeur de <tt>i</tt>.
| Calcule séquentiellement de <tt>i=0</tt> à <tt>i=N-1</tt>, inclusivement. || Chaque unité de calcul exécute la fonction pour une seule valeur de <tt>i</tt>.
|}
|}


== La directive <tt>kernels</tt> ==
== La directive <code>kernels</code> ==
Cette directive est dite ''descriptive''.  Le programmeur l'utilise pour signifier au compilateur les portions qui selon lui peuvent être parallélisées. Le compilateur fait ce qu'il veut de cette information et adopte la stratégie qui lui semble la meilleure pour exécuter le code, '''incluant''' son exécution séquentielle. De façon générale, le compilateur  
Cette directive est dite ''descriptive''.  Le programmeur l'utilise pour signifier au compilateur les portions qui selon lui peuvent être parallélisées. Le compilateur fait ce qu'il veut de cette information et adopte la stratégie qui lui semble la meilleure pour exécuter le code, '''incluant''' son exécution séquentielle. De façon générale, le compilateur  
# analyse le code pour détecter le parallélisme,
# analyse le code pour détecter le parallélisme,
Line 86: Line 102:
#pragma acc kernels
#pragma acc kernels
{
{
for (int i=0; i<N; i++)
  for (int i=0; i<N; i++)
{
  {
  C[i] = A[i] + B[i];
    C[i] = A[i] + B[i];
}
  }
}
}
</syntaxhighlight>  
</syntaxhighlight>  


Il est rare que le code soit aussi simple et il faut se baser sur la rétroaction du compilateur pour trouver les portions qu'il a négligé de paralléliser.  
Il est rare que le code soit aussi simple et il faut se baser sur la [[OpenACC_Tutorial_-_Profiling/fr#Renseignements_sur_le_compilateur| rétroaction du compilateur]] pour trouver les portions qu'il a négligé de paralléliser.  
 
{{Callout
{{Callout
|title=Description ou prescription
|title=Description ou prescription
|content=
|content=
Les utilisateurs d'[https://docs.computecanada.ca/wiki/OpenMP/fr OpenMP] retrouveront dans OpenACC le principe de directives. Il existe cependant d'importantes différences entre les directives OpenMP et OpenACC. Les directives OpenMP sont à la base ''prescriptives''. Ceci signifie que le compilateur est forcé d'accomplir la parallélisation, peu importe que l'effet détériore ou améliore la performance. Le résultat est prévisible pour tous les compilateurs. De plus, la parallélisation se fera de la même manière, peu importe le matériel utilisé pour exécuter le code. Par contre, le même code peut connaitre une moins bonne performance, dépendant de l'architecture. Il peut donc être préférable par exemple de changer l'ordre des boucles. Pour paralléliser du code avec OpenMP et obtenir une performance optimale dans différentes architectures, il faudrait avoir un ensemble différent de directives pour chaque architecture.  
Si vous avez déjà utilisé [[OpenMP/fr|OpenMP]], vous retrouverez dans OpenACC le principe de ''directives''. Il existe cependant d'importantes différences entre les directives OpenMP et OpenACC&nbsp;:
* Les directives OpenMP sont à la base ''prescriptives''. Ceci signifie que le compilateur est forcé d'accomplir la parallélisation, peu importe que l'effet détériore ou améliore la performance. Le résultat est prévisible pour tous les compilateurs. De plus, la parallélisation se fera de la même manière, peu importe le matériel utilisé pour exécuter le code. Par contre, le même code peut connaitre une moins bonne performance, dépendant de l'architecture. Il peut donc être préférable par exemple de changer l'ordre des boucles. Pour paralléliser du code avec OpenMP et obtenir une performance optimale dans différentes architectures, il faudrait avoir un ensemble différent de directives pour chaque architecture.  


Pour leur part, plusieurs directives OpenACC sont de nature ''descriptive''. Ici, le compilateur est libre de compiler le code de la façon qu'il juge la meilleure, selon l'architecture visée. Dans certains cas, le code ne sera pas parallélisé du tout. Le '''même code''' exécuté sur un GPU, sur un Xeon Phi ou sur un CPU peut donner du code binaire différent. Ceci signifie que la performance pourrait varier selon le compilateur  et que les compilateurs d'une nouvelle génération seront plus efficaces, surtout en présence de nouveau matériel.
* Pour leur part, plusieurs directives OpenACC sont de nature ''descriptive''. Ici, le compilateur est libre de compiler le code de la façon qu'il juge la meilleure, selon l'architecture visée. Dans certains cas, le code ne sera pas parallélisé du tout. Le '''même code''' exécuté sur un GPU ou sur un CPU peut donner du code binaire différent. Ceci signifie que la performance pourrait varier selon le compilateur  et que les compilateurs d'une nouvelle génération seront plus efficaces, surtout en présence de nouveau matériel.
}}
}}


=== Exemple : porter un produit matrice-vecteur ===  
=== Exemple : porter un produit matrice-vecteur ===  
Pour notre exemple, nous utilisons du code provenant du  [https://github.com/calculquebec/cq-formation-openacc répertoire Github], particulièrement une portion de code du fichier <tt>matrix_functions.h</tt>. Le code Fortran équivalent se trouve dans la sous-routine <tt>matvec</tt> contenue dans le fichier <tt>matrix.F90</tt>. Le code original est comme suit&nbsp;:
Pour notre exemple, nous utilisons du code provenant du  [https://github.com/calculquebec/cq-formation-openacc répertoire Github], particulièrement une portion de code [https://github.com/calculquebec/cq-formation-openacc/blob/main/cpp/matrix_functions.h#L20 fichier <code>cpp/matrix_functions.h</code>]. Le code Fortran équivalent se trouve dans la sous-routine [https://github.com/calculquebec/cq-formation-openacc/blob/main/f90/matrix.F90#L101 <code>matvec</code> contenue dans le fichier <code>matrix.F90</code>]. Le code C++ est comme suit&nbsp;:
<syntaxhighlight lang="cpp" line>
<syntaxhighlight lang="cpp" line start="29">
for(int i=0;i<num_rows;i++) {
  for(int i=0;i<num_rows;i++) {
  double sum=0;
    double sum=0;
  int row_start=row_offsets[i];
    int row_start=row_offsets[i];
  int row_end=row_offsets[i+1];
    int row_end=row_offsets[i+1];
  for(int j=row_start;j<row_end;j++) {
    for(int j=row_start;j<row_end;j++) {
    unsigned int Acol=cols[j];
      unsigned int Acol=cols[j];
    double Acoef=Acoefs[j];
      double Acoef=Acoefs[j];
    double xcoef=xcoefs[Acol];
      double xcoef=xcoefs[Acol];
    sum+=Acoef*xcoef;
      sum+=Acoef*xcoef;
    }
    ycoefs[i]=sum;
   }
   }
  ycoefs[i]=sum;
}
</syntaxhighlight>  
</syntaxhighlight>  


Le premier changement à faire au code est d'ajouter la directive <tt>kernels</tt>  pour essayer de le faire exécuter sur le GPU. Pour l'instant, nous n'avons pas à nous préoccuper du transfert des données ou à fournir des renseignements au compilateur.
Le [https://github.com/calculquebec/cq-formation-openacc/blob/main/cpp/step1.kernels/matrix_functions.h#L29  premier changement] à faire au code est d'ajouter la directive <tt>kernels</tt>  pour essayer de le faire exécuter sur le GPU.
<syntaxhighlight lang="cpp" line highlight="1,2,15">
Pour l'instant, nous n'avons pas à nous préoccuper du transfert des données ou à fournir des renseignements au compilateur.
 
<syntaxhighlight lang="cpp" line start="29" highlight="1,2,15">
#pragma acc kernels
#pragma acc kernels
   {
   {
Line 140: Line 160:
==== Construire avec OpenACC ====
==== Construire avec OpenACC ====


{{Callout
Les compilateurs NVidia utilisent l'option <code>-acc</code> pour permettre la compilation pour un accélérateur. Nous utilisons la sous-option <code>-gpu=managed</code> pour indiquer au compilateur que nous voulons utiliser la [https://developer.nvidia.com/blog/unified-memory-cuda-beginners/ mémoire gérée]  pour simplifier le transfert de données en provenance et à destination du périphérique; nous n'utiliserons pas cette option dans un prochain exemple. Nous utilisons aussi l'option <code>-fast</code> pour l'optimisation.
|title=
Choix du compilateur
|content=
En date de mai 2021, plusieurs compilateurs offraient les fonctionnalités OpenACC. Puisque cette technologie est principalement soutenue par [http://www.nvidia.com/content/global/global.php NVidia] et [http://www.cray.com/ Cray], les compilateurs de ces deux compagnies offrent le meilleur support pour OpenACC. Les fonctionnalités OpenACC des compilateurs de [https://gcc.gnu.org/wiki/OpenACC GNU] s'améliorent depuis la version 5.
 
Dans ce tutoriel, nous utilisons la version 20.7 des compilateurs NVidia pour le calcul de haute performance.}}
 
Les compilateurs NVidia utilisent l'option <tt>-ta</tt> (''target accelerator'') pour permettre la compilation pour un accélérateur. Nous utilisons la sous-option <tt>tesla:managed</tt> pour indiquer au compilateur que nous voulons compiler pour des GPU Tesla et que nous voulons utiliser la mémoire gérée pour simplifier le transfert de données en provenance et à destination du périphérique; nous n'utiliserons pas cette option dans un prochain exemple. Nous utilisons aussi l'option <tt>-fast</tt> pour l'optimisation.


{{Command
{{Command
|nvc++ -fast -Minfo{{=}}accel -ta{{=}}tesla:managed main.cpp -o challenge
|nvc++ -fast -Minfo{{=}}accel -acc -gpu{{=}}managed main.cpp -o challenge
|result=
|result=
...
...
matvec(const matrix &, const vector &, const vector &):
matvec(const matrix &, const vector &, const vector &):
     23, include "matrix_functions.h"
     23, include "matrix_functions.h"
           27, Generating implicit copyin(xcoefs[:]) [if not already present]
           30, Generating implicit copyin(cols[:],row_offsets[:num_rows+1],Acoefs[:]) [if not already present]
               Generating implicit copyout(ycoefs[:num_rows]) [if not already present]
               Generating implicit copyout(ycoefs[:num_rows]) [if not already present]
               Generating implicit copyin(row_offsets[:num_rows+1],Acoefs[:],cols[:]) [if not already present]
               Generating implicit copyin(xcoefs[:]) [if not already present]
           30, Loop carried dependence of ycoefs-> prevents parallelization
           31, Loop carried dependence of ycoefs-> prevents parallelization
               Loop carried backward dependence of ycoefs-> prevents vectorization
               Loop carried backward dependence of ycoefs-> prevents vectorization
               Complex loop carried dependence of Acoefs->,xcoefs-> prevents parallelization
               Complex loop carried dependence of Acoefs->,xcoefs-> prevents parallelization
               Generating Tesla code
               Generating NVIDIA GPU code
               30, #pragma acc loop seq
               31, #pragma acc loop seq
               34, #pragma acc loop vector(128) /* threadIdx.x */
               35, #pragma acc loop vector(128) /* threadIdx.x */
                   Generating implicit reduction(+:sum)
                   Generating implicit reduction(+:sum)
           34, Loop is parallelizable
           35, Loop is parallelizable
}}
}}


Le résultat montre que les deux boucles n'ont pas pu être parallélisées par le compilateur; ci-dessous, nous expliquons comment éviter ceci.  
Le résultat montre que la boucle externe sur la ligne 31 n'a pas pu être parallélisée par le compilateur. Dans la prochaine section, nous expliquons comment traiter ces dépendances.


== Réparer les fausses dépendances de boucles ==
== Réparer les fausses dépendances de boucles ==
Line 186: Line 198:


=== Boucle avec clause <tt>independent</tt> ===  
=== Boucle avec clause <tt>independent</tt> ===  
Une autre façon de s'assurer que le compilateur traite les boucles de manière indépendante est de le spécifier explicitement  avec la clause <tt>independent</tt>. Comme toute autre directive ''prescriptive'', le compilateur y est obligé et l'analyse qu'il pourrait faire ne sera pas considérée. En reprenant l'exemple de la section ''La directive <tt>kernels</tt>'', nous avons&nbsp;:   
Une autre façon de s'assurer que le compilateur traite les boucles de manière indépendante est de le spécifier explicitement  avec la clause <tt>independent</tt>. Comme toute autre directive ''prescriptive'', le compilateur y est obligé et l'analyse qu'il pourrait faire ne sera pas considérée. En reprenant l'exemple de la section ''La directive <tt>kernels</tt>'' ci-dessus, nous avons&nbsp;:   
<syntaxhighlight lang="cpp" line highlight="3">
<syntaxhighlight lang="cpp" line highlight="3">
#pragma acc kernels
#pragma acc kernels
Line 213: Line 225:
Remarquez que les autres pointeurs n'ont pas besoin d'être restreints puisque le compilateur ne les rapporte pas comme causant des problèmes. En recompilant avec les changements que nous venons de faire, le compilateur émet le message suivant&nbsp;:  
Remarquez que les autres pointeurs n'ont pas besoin d'être restreints puisque le compilateur ne les rapporte pas comme causant des problèmes. En recompilant avec les changements que nous venons de faire, le compilateur émet le message suivant&nbsp;:  
{{Command
{{Command
|nvc++ -fast -Minfo{{=}}accel -ta{{=}}tesla:managed main.cpp -o challenge
|nvc++ -fast -Minfo{{=}}accel -acc -gpu{{=}}managed main.cpp -o challenge
|result=
|result=
matvec(const matrix &, const vector &, const vector &):
matvec(const matrix &, const vector &, const vector &):
Line 228: Line 240:


== Performance du code porté ==  
== Performance du code porté ==  
Maintenant que le code est porté sur le GPU, nous pouvons analyser sa performance et vérifier si les résultats sont corrects. L'exécution du code original sur un des nœuds GPU de Guillimin produit ceci&nbsp;:
Maintenant que le code est porté sur le GPU, nous pouvons analyser sa performance et vérifier si les résultats sont corrects. L'exécution du code original sur un nœud GPU produit ceci&nbsp;:
{{Command
{{Command
|./cg.x  
|./cg.x  
Line 263: Line 275:
}}
}}
[[File:Openacc profiling1.png|thumbnail|Cliquez pour agrandir.]]
[[File:Openacc profiling1.png|thumbnail|Cliquez pour agrandir.]]
Les résultats sont corrects,  toutefois, loin de gagner en vitesse, l'opération a pris près de quatre fois plus de temps! Utilisons le NVidia Visual Profiler (<tt>nvvp</tt>) pour voir ce qui se passe.  
Les résultats sont corrects,  toutefois, loin de gagner en vitesse, l'opération a pris près de quatre fois plus de temps! Utilisons le NVidia Visual Profiler (<tt>nvvp</tt>) pour voir ce qui se passe.
# Démarrez  <tt>nvvp</tt> avec la commande <tt>nvvp &</tt>, où le symbole <tt>&</tt> permet de démarrer en arrière-plan.
 
# Sélectionnez ''File -> New Session''.
=== NVIDIA Visual Profiler ===
# Dans le champ "File", cherchez l'exécutable; dans notre exemple, nous utilisons  <tt>challenge</tt>.
[[File:Nvvp-pic0.png|thumbnail|300px|Profileur NVVP|right]]
# Cliquez sur "Next" jusqu'à ce que vous puissiez cliquer sur "Finish".  
[[File:Nvvp-pic1.png|thumbnail|300px|Recherche de l'exécutable que vous voulez profiler|right]]
 
[https://developer.nvidia.com/nvidia-visual-profiler NVIDIA Visual Profiler (NVVP)] est un profileur graphique pour les applications OpenACC.
C'est un outil d'analyse pour les '''codes écrits avec les directives OpenACC et CUDA C/C++'''.
En conséquence, si l'exécutable n'utilise pas le GPU, ce profileur ne fournira aucun résultat.
 
Quand [[Visualization#Fenêtres_à_distance_avec_redirection_X11|X11 est redirigé vers un serveur X-Server]] ou quand vous utilisez un [[VNC/fr|environnement bureau Linux]] (aussi via [[JupyterHub/fr#Bureau|JupyterHub]] avec 2 cœurs CPU, 5000M de mémoire et 1 GPU),
vous pouvez lancer NVVP à partir d'un terminal&nbsp;:
{{Command
|module load cuda/11.7 java/1.8
}}
{{Command
|nvvp
}}


Le programme est exécuté et on obtient un tableau chronologique du déroulement (voir l'image). On remarque que le transfert de données entre le départ et l'arrivée occupe la plus grande partie du temps d'exécution, ce qui est fréquent quand du code est porté d'un CPU vers un GPU. Nous verrons comment ceci peut être amélioré dans la prochaine partie, [https://docs.computecanada.ca/wiki/OpenACC_Tutorial_-_Data_movement/fr Mouvement des données].
# Après l'affichage de la fenêtre de lancement de NVVP, vous devez entrer le répertoire ''Workspace'' qui sera employé pour les fichiers temporaires. Dans le chemin suggéré, remplacez <code>home</code> par <code>scratch</code> et cliquez sur ''OK''.
# Sélectionnez ''File > New Session'' ou cliquez sur le bouton correspondant dans la barre d'outils.
# Cliquez sur le bouton ''Browse'' à la droite du champ ''File'' pour le chemin.
## Changez le répertoire s'il y a lieu.
## Sélectionnez un exécutable construit avec des codes écrits avec des directives OpenACC et CUDA C/C++.
# Sous le champ ''Arguments'', sélectionnez l'option ''Profile current process only''.
# Cliquez sur ''Next >'' pour voir les autres options de profilage.
# Cliquez sur ''Finish'' pour lancer le profilage de l'exécutable.
 
Pour faire ceci, suivez ces étapes&nbsp;:
# Lancez <tt>nvvp</tt> avec la commande <tt>nvvp &</tt>  (le symbole <tt>&</tt> commande le lancement en arrière-plan).
# Sélectionnez '' File -> New Session''.
# Dans le champ ''File:'', cherchez l'exécutable (nommé dans notre exemple <tt>challenge</tt>).
# Cliquez sur ''Next'' jusqu'à ce que vous puissiez cliquer sur ''Finish''.
 
Le programme est exécuté et on obtient un tableau chronologique du déroulement (voir l'image). On remarque que le transfert de données entre le départ et l'arrivée occupe la plus grande partie du temps d'exécution, ce qui est fréquent quand du code est porté d'un CPU vers un GPU. Nous verrons comment ceci peut être amélioré dans la prochaine partie, [[OpenACC Tutorial - Data movement/fr|Mouvement des données]].


== La directive <tt>parallel loop</tt> ==  
== La directive <tt>parallel loop</tt> ==  
Line 306: Line 346:
La compilation produit le message suivant&nbsp;:
La compilation produit le message suivant&nbsp;:
{{Command
{{Command
|pgc++ -fast -Minfo{{=}}accel -ta{{=}}tesla:managed main.cpp -o challenge
|nvc++ -fast -Minfo{{=}}accel -acc -gpu{{=}}managed main.cpp -o challenge
|result=
|result=
matvec(const matrix &, const vector &, const vector &):
matvec(const matrix &, const vector &, const vector &):
Line 319: Line 359:
           34, Loop is parallelizable
           34, Loop is parallelizable
}}
}}
==Différences entre <tt>parallel loop</tt> et <tt>kernels</tt>==
==Différences entre <tt>parallel loop</tt> et <tt>kernels</tt>==
{| class="wikitable" width="100%"
{| class="wikitable" width="100%"
Line 341: Line 382:
|content=
|content=
# Modifiez les fonctions <tt>matvec</tt>, <tt>waxpby</tt> et <tt>dot</tt>. Vous pouvez utiliser soit <tt>kernels</tt>, soit <tt>parallel loop</tt>. La solution se trouve dans les répertoires <tt>step1. de [https://github.com/calculquebec/cq-formation-openacc Github] *</tt>.  
# Modifiez les fonctions <tt>matvec</tt>, <tt>waxpby</tt> et <tt>dot</tt>. Vous pouvez utiliser soit <tt>kernels</tt>, soit <tt>parallel loop</tt>. La solution se trouve dans les répertoires <tt>step1. de [https://github.com/calculquebec/cq-formation-openacc Github] *</tt>.  
# Modifiez le Makefile en ajoutant <tt>-ta{{=}}tesla:managed</tt> et <tt>-Minfo{{=}}accel</tt> aux indicateurs pour le compilateur.  
# Modifiez le Makefile en ajoutant <tt>-acc -gpu{{=}}managed</tt> et <tt>-Minfo{{=}}accel</tt> aux indicateurs pour le compilateur.  
}}
}}


[https://docs.computecanada.ca/wiki/OpenACC_Tutorial_-_Data_movement/fr Page suivante, Mouvement des données]<br>
[[OpenACC Tutorial - Profiling/fr|<- Page précédente, Profileurs]] | [[OpenACC Tutorial/fr|^- Retour au début du tutoriel]] | [[OpenACC Tutorial - Data movement/fr|Page suivante, Mouvement des données ->]]
[https://docs.computecanada.ca/wiki/OpenACC_Tutorial/fr Retour au début du tutoriel]
FuzzyBot
Bots
38,760

edits

Retrieved from "https://docs.alliancecan.ca/wiki/Special:MobileDiff/101513...138004"