OpenACC Tutorial - Adding directives/fr: Difference between revisions

Jump to navigation Jump to search
Updating to match new version of source page
No edit summary
(Updating to match new version of source page)
Line 96: Line 96:
|}
|}


<div class="mw-translate-fuzzy">
== La directive <tt>kernels</tt> ==
== La directive <tt>kernels</tt> ==
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  
Line 102: Line 103:
# crée un kernel,
# crée un kernel,
# transfère le kernel au GPU.
# transfère le kernel au GPU.
</div>


Voici un exemple de cette directive&nbsp;:
Voici un exemple de cette directive&nbsp;:
Line 108: Line 110:
#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.  
<div class="mw-translate-fuzzy">
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.
</div>
 
{{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.  
<div class="mw-translate-fuzzy">
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.
</div>


<div class="mw-translate-fuzzy">
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, 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.
</div>
}}
}}


<div class="mw-translate-fuzzy">
=== 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 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;:
<syntaxhighlight lang="cpp" line>
</div>
for(int i=0;i<num_rows;i++) {
<syntaxhighlight lang="cpp" line start="29">
  double sum=0;
  for(int i=0;i<num_rows;i++) {
  int row_start=row_offsets[i];
    double sum=0;
  int row_end=row_offsets[i+1];
    int row_start=row_offsets[i];
  for(int j=row_start;j<row_end;j++) {
    int row_end=row_offsets[i+1];
    unsigned int Acol=cols[j];
    for(int j=row_start;j<row_end;j++) {
    double Acoef=Acoefs[j];
      unsigned int Acol=cols[j];
    double xcoef=xcoefs[Acol];
      double Acoef=Acoefs[j];
    sum+=Acoef*xcoef;
      double xcoef=xcoefs[Acol];
      sum+=Acoef*xcoef;
    }
    ycoefs[i]=sum;
   }
   }
  ycoefs[i]=sum;
}
</syntaxhighlight>  
</syntaxhighlight>  


<div class="mw-translate-fuzzy">
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 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.
<syntaxhighlight lang="cpp" line highlight="1,2,15">
</div>
 
<syntaxhighlight lang="cpp" line start="29" highlight="1,2,15">
#pragma acc kernels
#pragma acc kernels
   {
   {
Line 162: Line 176:
==== Construire avec OpenACC ====
==== Construire avec OpenACC ====


{{Callout
<div class="mw-translate-fuzzy">
|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.
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.
</div>


{{Command
{{Command
Line 178: Line 186:
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.  
<div class="mw-translate-fuzzy">
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.
</div>


== Réparer les fausses dépendances de boucles ==
== Réparer les fausses dépendances de boucles ==
38,760

edits

Navigation menu