OpenACC Tutorial - Adding directives/fr: Difference between revisions

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''.

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.  

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.  

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.

=== Boucles et kernels ===

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.

void loopBody(A,B,C,i)

|Calcule de 0-N, dans l'ordre || Chaque cœur de calcul traite une valeur de <tt>i</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équencielle. De façon générale, le compilateur  

for (int i=0; i<N; i++)

{

  C[i] = A[i] + B[i];

}

Il est rare que du 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.  

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.  

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 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 sousroutine <tt>matvec</tt> contenue dans le fichier <tt>matrix.F90</tt>. Le code original est comme suit&nbsp;:

<syntaxhighlight lang="cpp" line>

for(int i=0;i<num_rows;i++) {

  double sum=0;

  int row_start=row_offsets[i];

  int row_end=row_offsets[i+1];

  for(int j=row_start;j<row_end;j++) {

    unsigned int Acol=cols[j];

    double Acoef=Acoefs[j];

    double xcoef=xcoefs[Acol];

    sum+=Acoef*xcoef;

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">

==== Développer avec OpenACC ====  

Dans ce tutoriel, nous travaillons avec la version 16.3 des [http://www.pgroup.com/support/download_pgi2016.php?view=current compilateurs du Portland Group]. Nous utilisons l'option <tt>-ta</tt> (pour ''target accelerator'') pour porter le code sur les accélérateurs et la sous-option <tt>tesla:managed</tt> pour signifier au compilateur que les GPU Tesla sont la cible et que nous voulons utiliser une mémoire autogérée, ce qui simplifie le transfert de données dans les deux directions. Cette option ne sera pas utilisée dans un prochain exemple. Nous utilisons aussi l'option d'optimisation <tt>-fast</tt>.

En résultat, nous constatons que le compilateur n'a pu paralléliser les deux boucles;  nous verrons plus loin comment traiter ce cas.

|pgc++ -fast -Minfo{{=}}accel -ta{{=}}tesla:managed main.cpp -o challenge

           27, Generating copyout(ycoefs[:num_rows])

               Generating copyin(xcoefs[:],Acoefs[:],cols[:],row_offsets[:num_rows+1])

           29, Complex loop carried dependence of row_offsets-> prevents parallelization

              Loop carried dependence of ycoefs-> prevents parallelization

               Complex loop carried dependence of cols->,Acoefs->,xcoefs-> prevents parallelization

               Accelerator kernel generated

               Generating Tesla code

               33, #pragma acc loop vector(128) /* threadIdx.x */

              37, Sum reduction generated for sum

           33, Loop is parallelizable

 

Même lorsque le programmeur sait qu'une boucle peut être parallélisée, il arrive que le compilateur ne le remarque pas. Un cas commun en C/C++ est connu sous le nom de '[https://en.wikipedia.org/wiki/Pointer_aliasing ''pointer aliasing'']. Contrairement au Fortran, C/C++ ne possèdent pas comme tel de tableaux (''arrays''), mais plutôt des pointeurs. Le concept d'alias s'applique à deux pointeurs dirigés vers la même mémoire. Si le compilateur ne sait pas que des pointeurs ne sont pas des alias, il doit cependant le supposer. Dans l'exemple précédent, on voit clairement pourquoi le compilateur ne pouvait pas paralléliser la boucle. En supposant que les pointeurs sont identiques, il y a forcément dépendance des itérations de la boucle.  

Une des manières de dire au compilateur que les pointeurs '''ne sont pas''' des alias est d'utiliser  le mot-clé  <tt>restrict</tt>, introduit à cette fin dans C99.  Il n'y a toujours pas de manière standard pour ce faire en C++, mais chaque compilateur possède un mot-clé qui lui est propre. Dépendant du compilateur, on peut utiliser <tt>__restrict</tt> ou <tt>__restrict__</tt>. Les compilateurs du Portland Group utilisent <tt>__restrict</tt>. Pour savoir pourquoi il n'existe pas de standard en C++, consultez [http://www.open-std.org/jtc1/sc22/wg21/docs/papers/2014/n3988.pdf ce document]. Ce concept est important pour OpenACC comme pour toute programmation C/C++, car les compilateurs peuvent effectuer plusieurs autres optimisations si les pointeurs ne sont pas des alias. Remarquez que le mot-clé se place '''après''' le pointeur puisque c'est à ce dernier qu'il se réfère, et non au type; autrement dit, la déclaration doit se lire <code>float * __restrict A;</code> plutôt que <code>float __restrict * A;</code>.  

En déclarant un pointeur comme étant ''restraint'', on s'assure qu'uniquement ce pointeur ou une valeur dérivée (comme <tt>ptr +1</tt>) pourra accéder à l'objet auquel il réfère, et ce pour la durée de vie du pointeur. Ceci est une garantie que le programmeur donne au compilateur;  si le programmeur manque à son obligation, le comportement n'est pas défini. Pour plus d'information, consultez l'article Wikipédia [https://en.wikipedia.org/wiki/Fstab restrict].  

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;: :

Revenons au cas du produit matrice-vecteur présenté plus haut. Notre recommandation pour éviter les faux alias est de définir les pointeurs comme étant restreints en remplaçant le code de <tt>matrix_functions.h</tt>.  

|pgc++ -fast -Minfo{{=}}accel -ta{{=}}tesla:managed main.cpp -o challenge

           27, Generating copyout(ycoefs[:num_rows])

               Generating copyin(xcoefs[:],Acoefs[:],cols[:],row_offsets[:num_rows+1])

           29, Loop is parallelizable

               29, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */

           33, Loop is parallelizable

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;:

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.   

# 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, [https://docs.computecanada.ca/wiki/OpenACC_Tutorial_-_Data_movement/fr Mouvement des données].

Avec la directive <tt>kernels</tt>, c'est le compilateur qui fait toute l'analyse; ceci est une approche ''descriptive'' pour porter du code. OpenACC offre aussi une approche ''prescriptive'' avec la directive <tt>parallel</tt> qui peut être combinée à la directive <tt>loop</tt> ainsi&nbsp;:

* Avec la clause <tt>reduction</tt> , les valeurs de la variable dans chaque itération est ''réduite'' à une valeur unique.  La clause s'utilise entre autres avec les opérations addition (+), multiplication (*), maximum (max) et minimum (min).  

Compiling this code yields the following compiler feedback:

|pgc++ -fast -Minfo{{=}}accel -ta{{=}}tesla:managed main.cpp -o challenge

== Parallel loop vs kernel ==

! PARALLEL LOOP  !! KERNEL

* It is the programmer's responsibility to ensure that parallelism is safe

* Enables parallelization of sections that the compiler may miss

* Straightforward path from OpenMP

* It is the compiler's responsibility to analyze the code and determine what is safe to parallelize.

* A single directive can cover a large area of code

* The compiler has more room to optimize

Both approaches are equally valid and can perform equally well.

  {{Challenge

|title=Challenge: Add OpenACC directives <tt>kernels</tt> or <tt>parallel loop</tt>  

# Modify the functions <tt>matvec</tt>, <tt>waxpby</tt> and <tt>dot</tt> to use OpenACC. You may use either the <tt>kernels</tt> or the <tt>parallel loop</tt> directives. The directories <tt>step1.*</tt> contain the solution.  

# Modify the Makefile to add <tt>-ta{{=}}tesla:managed</tt> and <tt>-Minfo{{=}}accel</tt> to your compiler flags.  

[[OpenACC Tutorial - Data movement|Onward to the next unit: Data movement]]<br>

[[OpenACC Tutorial|Back to the lesson plan]]