View source for OpenACC Tutorial - Adding directives/fr

<languages />

{{Objectives
|title=Objectifs d'apprentissage
|content=
* comprendre le processus de transfert (''offloading'')
* comprendre ce qu'est une directive OpenACC 
* connaitre la différence entre les directives <tt>loop</tt> et <tt>kernels</tt>  
* savoir programmer avec OpenACC
* comprendre le concept d'alias en C/C++
* savoir utiliser la rétroaction du compilateur et éviter les faux alias
}}

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

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

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%"
|-
! C/C++ !! FORTRAN
|-
| <syntaxhighlight lang="cpp" line highlight="1,2,13">
#pragma acc kernels
{
  for (int i=0; i<N; i++)
  {
    x[i] = 1.0;
    y[i] = 2.0;
  }

  for (int i=0; i<N; i++)
  {
    y[i] = a * x[i] + y[i];
  }
}
</syntaxhighlight> || <syntaxhighlight lang="fortran" line highlight="1,7">
!$acc kernels
  do i=1,N
    x(i) = 1.0
    y(i) = 2.0
  end do
  
  y(:) = a*x(:) + y(:)
!$acc end kernels
</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.

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

=== 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.
{| class="wikitable" width="100%"
|-
! BOUCLE !! KERNEL
|-
| <syntaxhighlight lang="cpp" line>
for (int i=0; i<N; i++)
{
  C[i] = A[i] + B[i];
}
</syntaxhighlight> || <syntaxhighlight lang="cpp" line>
void loopBody(A,B,C,i)
{
  C[i] = A[i] + B[i];
}
</syntaxhighlight>
|-
|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 
# analyse le code pour détecter le parallélisme,
# s'il détecte du parallélisme, identifie les données à transférer et décide quand faire le transfert,
# crée un kernel,
# transfère le kernel au GPU.

Voici un exemple de cette directive&nbsp;:

<syntaxhighlight lang="cpp" line highlight="1,2,7">
#pragma acc kernels
{
for (int i=0; i<N; i++)
{
  C[i] = A[i] + B[i];
}
}
</syntaxhighlight> 

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. 
{{Callout
|title=Description ou prescription
|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. 

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

=== 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 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;
  }
  ycoefs[i]=sum;
}
</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.
<syntaxhighlight lang="cpp" line highlight="1,2,15">
#pragma acc kernels
  {
    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;
      }
      ycoefs[i]=sum;
    }
  }
</syntaxhighlight> 

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

{{Command
|pgc++ -fast -Minfo{{=}}accel -ta{{=}}tesla:managed main.cpp -o challenge
|result=
...
matvec(const matrix &, const vector &, const vector &):
     23, include "matrix_functions.h"
          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
              Loop carried backward dependence of ycoefs-> prevents vectorization
              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
}}
As we can see in the compiler output, the compiler could not parallelize the two loops. We will see in the following sections how to deal with this. 
{{Callout
|title=Choix du compilateur
|content=
En date de mai 2016, relativement peu de compilateurs offraient les fonctionnalités d'OpenACC. Les plus avancés en ce sens sont les compilateurs du [http://www.pgroup.com/ Portland Group] de [http://www.nvidia.com/content/global/global.php NVidia] et ceux de [http://www.cray.com/ Cray]. Pour ce qui de [https://gcc.gnu.org/wiki/OpenACC GNU], l'implémentation d'OpenACC était expérimentale et devrait être complète dans la version 6. 

Dans ce tutoriel, nous utilisons la version 16.3 des [http://www.pgroup.com/support/download_pgi2016.php?view=current compilateurs du Portland Group]  qui sont gratuits pour des fins de recherche universitaire. 
}}


== Réparer les fausses dépendances de boucles ==
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. 

=== <tt>restrict</tt> keyword ===
One way to tell the compiler that pointers are '''not''' going to be aliased, is by using a special keyword. In C, the keyword <tt>restrict</tt> was introduced in C99 for this purpose.  In C++, there is no standard way yet, but each compiler typically has its own keyword. Either <tt>__restrict</tt> or <tt>__restrict__</tt> can be used depending on the compiler. For Portland Group compilers, the keyword is <tt>__restrict</tt>. For an explanation as to why there is no standard way to do this in C++, you can read [http://www.open-std.org/jtc1/sc22/wg21/docs/papers/2014/n3988.pdf this paper]. This concept is important not only for OpenACC, but for any C/C++ programming, since many more optimizations can be done by compilers when pointers are guaranteed not to be aliased. Note that the keyword goes ''after'' the pointer, since it refers to the pointer, and not to the type. In other words, you would declare <code>float * __restrict A;</code> rather than <code>float __restrict * A;</code>. 


{{Callout
|title=What does <tt>restrict</tt> really mean ?
|content=
Declaring a pointer as restricted formally means that for "the lifetime of the pointer, only it or a value derived from it (such as <tt>ptr +1</tt>) will be used to access the object to which it points". This is a guarantee that the ''programmer'' gives to the ''compiler''. If the programmer violates this guarantee, behaviour is undefined. For more information on this concept, see this [https://en.wikipedia.org/wiki/Restrict Wikipedia article]. 
}}

=== Loop directive with independent clause ===
Another way to tell the compiler that loops iterations are independent is to specify it explicitly by using a different directive: <tt>loop</tt>, with the clause <tt>independent</tt>. This is a ''prescriptive'' directive. Like any prescriptive directive, this tells the compiler what to do, and overrides any compiler analysis. The initial example above would become: 
<syntaxhighlight lang="cpp" line highlight="3">
#pragma acc kernels
{
#pragma acc loop independent
for (int i=0; i<N; i++)
{
  C[i] = A[i] + B[i];
}
}
</syntaxhighlight> 

=== Back to the example ===
Going back to the matrix-vector product above, the way that we recommend fixing false aliasing is by declaring the pointers as restricted. This is done by changing the following code in <tt>matrix_functions.h</tt>: 
<syntaxhighlight lang="cpp" line>
  double *Acoefs=A.coefs;
  double *xcoefs=x.coefs;
  double *ycoefs=y.coefs;
</syntaxhighlight> 
by this code: 
<syntaxhighlight lang="cpp" line>
  double *__restrict Acoefs=A.coefs;
  double *__restrict xcoefs=x.coefs;
  double *__restrict ycoefs=y.coefs;
</syntaxhighlight> 
We note that we do not need to declare the other pointers as restricted, since they are not reported as problematic by the compiler. With the above changes, recompiling gives the following compiler messages: 
{{Command
|pgc++ -fast -Minfo{{=}}accel -ta{{=}}tesla:managed main.cpp -o challenge
|result=
matvec(const matrix &, const vector &, const vector &):
     23, include "matrix_functions.h"
          27, Generating copyout(ycoefs[:num_rows])
              Generating copyin(xcoefs[:],Acoefs[:],cols[:],row_offsets[:num_rows+1])
          29, Loop is parallelizable
              Accelerator kernel generated
              Generating Tesla code
              29, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
          33, Loop is parallelizable
}}

== How is ported code performing ? ==
Since we have completed a first step to porting the code to GPU, we need to analyze how the code is performing, and whether it gives the correct results. Running the original version of the code yields the following (performed on one of Guillimin's GPU node):
{{Command
|./cg.x 
|result=
Rows: 8120601, nnz: 218535025
Iteration: 0, Tolerance: 4.0067e+08
Iteration: 10, Tolerance: 1.8772e+07
Iteration: 20, Tolerance: 6.4359e+05
Iteration: 30, Tolerance: 2.3202e+04
Iteration: 40, Tolerance: 8.3565e+02
Iteration: 50, Tolerance: 3.0039e+01
Iteration: 60, Tolerance: 1.0764e+00
Iteration: 70, Tolerance: 3.8360e-02
Iteration: 80, Tolerance: 1.3515e-03
Iteration: 90, Tolerance: 4.6209e-05
Total Iterations: 100 Total Time: 29.894881s
}}

Running the OpenACC version yields the following:
{{Command
|./challenge 
|result=Rows: 8120601, nnz: 218535025
Iteration: 0, Tolerance: 4.0067e+08
Iteration: 10, Tolerance: 1.8772e+07
Iteration: 20, Tolerance: 6.4359e+05
Iteration: 30, Tolerance: 2.3202e+04
Iteration: 40, Tolerance: 8.3565e+02
Iteration: 50, Tolerance: 3.0039e+01
Iteration: 60, Tolerance: 1.0764e+00
Iteration: 70, Tolerance: 3.8360e-02
Iteration: 80, Tolerance: 1.3515e-03
Iteration: 90, Tolerance: 4.6209e-05
Total Iterations: 100 Total Time: 115.068931s
}}
[[File:Openacc profiling1.png|thumbnail|Click to enlarge]]
The results are correct. However, not only do we not get any speed up, but we rather get a slow down by a factor of almost 4! Let's profile the code again using NVidia's visual profiler (<tt>nvvp</tt>). This can be done with the following steps: 
# Start <tt>nvvp</tt> with the command <tt>nvvp &</tt>   (the <tt>&</tt> sign is to start it in the background)
# Go in File -> New Session
# In the "File:" field, search for the executable (named <tt>challenge</tt> in our example).
# Click "Next" until you can click "Finish". 

This will run the program and generate a timeline of the execution. The resulting timeline is illustrated on the image on the right side. As we can see, almost all of the run time is being spent transferring data between the host and the device. This is very often the case when one ports a code from CPU to GPU. We will look at how to optimize this in the [[OpenACC Tutorial - Data movement|next part of the tutorial]].

== The <tt>parallel loop</tt> directive ==
With the <tt>kernels</tt> directive, we let the compiler do all of the analysis. This is the ''descriptive'' approach to porting a code. OpenACC supports a ''prescriptive'' approach through a different directive, called the <tt>parallel</tt> directive. This can be combined with the <tt>loop</tt> directive, to form the <tt>parallel loop</tt> directive. An example would be the following code:
<syntaxhighlight lang="cpp" line highlight="1">
#pragma acc parallel loop
for (int i=0; i<N; i++)
{
  C[i] = A[i] + B[i];
}
</syntaxhighlight> 
Since <tt>parallel loop</tt> is a ''prescriptive'' directive, it forces the compiler to perform the loop in parallel. This means that the <tt>independent</tt> clause introduced above is implicit within a parallel region. 

For reasons that we explain below, in order to use this directive in the matrix-vector product example, we need to introduce additional clauses used to manage the scope of data. The <tt>private</tt> and <tt>reduction</tt> clauses control how the data flows through a parallel region. 
* With the <tt>private</tt> clause, a copy of the variable is made for each loop iteration, making the value of the variable independent from other iterations. 
* With the <tt>reduction</tt> clause, the values of a variable in each iteration will be ''reduced'' to a single value. It supports addition (+), multiplication (*), maximum (max), minimum (min), among other operations. 
These clauses were not required with the <tt>kernels</tt> directive, because the <tt>kernels</tt> directive handles this for you. 

Going back to the matrix-vector multiplication example, the corresponding code with the <tt>parallel loop</tt> directive would look like this:
<syntaxhighlight lang="cpp" line highlight="6">
#pragma acc parallel loop
  for(int i=0;i<num_rows;i++) {
    double sum=0;
    int row_start=row_offsets[i];
    int row_end=row_offsets[i+1];
#pragma acc loop reduction(+:sum)
    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;
    }
    ycoefs[i]=sum;
  }
</syntaxhighlight>
Compiling this code yields the following compiler feedback:
{{Command
|pgc++ -fast -Minfo{{=}}accel -ta{{=}}tesla:managed main.cpp -o challenge
|result=
matvec(const matrix &, const vector &, const vector &):
     23, include "matrix_functions.h"
          27, Accelerator kernel generated
              Generating Tesla code
              29, #pragma acc loop gang /* blockIdx.x */
              34, #pragma acc loop vector(128) /* threadIdx.x */
                  Sum reduction generated for sum
          27, Generating copyout(ycoefs[:num_rows])
              Generating copyin(xcoefs[:],Acoefs[:],cols[:],row_offsets[:num_rows+1])
          34, Loop is parallelizable
}}

== Parallel loop vs kernel ==
{| class="wikitable" width="100%"
|-
! 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> 
|content=
# 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]]