OpenACC Tutorial - Adding directives/fr

Objectifs d'apprentissage

• comprendre le processus de transfert (offloading)
• comprendre ce qu'est une directive OpenACC
• connaitre la différence entre les directives loop et kernels
• 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

Les directives OpenAcc sont semblables aux directives OpenMP. En C/C++, ce sont des pragmas 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 pragma à 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 pragmas 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 : la première initialise deux vecteurs et la seconde effectue une opération de niveau 1 d'addition des vecteurs.

C/C++	FORTRAN
#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]; } }	!$acc kernels do i=1,N x(i) = 1.0 y(i) = 2.0 end do y(:) = a*x(:) + y(:) !$acc end kernels

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

Loop et kernels

Quand le compilateur lit la directive OpenACC kernels, 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.

When the compiler reaches an OpenACC kernels directive, it will analyze the code in order to identify sections that can be parallelized. This often corresponds to the body of a loop that has independent iterations. When such a case is identified, the compiler will first wrap the body of the loop into a special function called a kernel. This internal code refactoring makes sure that each call to the kernel is independent from any other call. The kernel is then compiled to enable it to run on an accelerator. Since each call is independent, each one of the hundreds of cores of the accelerator can run the function for one specific index in parallel.

BOUCLE	KERNEL
for (int i=0; i<N; i++) { C[i] = A[i] + B[i]; }	void kernelName(A, B, C, i) { C[i] = A[i] + B[i]; }
Calcule séquentiellement de i=0 à i=N-1, inclusivement.	Chaque unité de calcul exécute la fonction pour une seule valeur de i.

La directive kernels

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

1. analyse le code pour détecter le parallélisme,
2. s'il détecte du parallélisme, identifie les données à transférer et décide quand faire le transfert,
3. crée un kernel,
4. transfère le kernel au GPU.

Voici un exemple de cette directive :

#pragma acc kernels
{
for (int i=0; i<N; i++)
{
  C[i] = A[i] + B[i];
}
}

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.

Description ou prescription

Les utilisateurs d'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 répertoire Github, particulièrement une portion de code du fichier matrix_functions.h. Le code Fortran équivalent se trouve dans la sous-routine matvec contenue dans le fichier matrix.F90. Le code original est comme suit :

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

Le premier changement à faire au code est d'ajouter la directive kernels 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.

#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;
    }
  }

Construire avec OpenACC

Choix du compilateur

En date de mai 2021, plusieurs compilateurs offraient les fonctionnalités OpenACC. Puisque cette technologie est principalement soutenue par NVidia et Cray, les compilateurs de ces deux compagnies offrent le meilleur support pour OpenACC. Les fonctionnalités OpenACC des compilateurs de 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 -ta (target accelerator) pour permettre la compilation pour un accélérateur. Nous utilisons la sous-option tesla:managed 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 -fast pour l'optimisation.

[name@server ~]$ nvc++ -fast -Minfo=accel -ta=tesla:managed main.cpp -o challenge
...
matvec(const matrix &, const vector &, const vector &):
     23, include "matrix_functions.h"
          27, Generating implicit copyin(xcoefs[:]) [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]
          30, Loop carried dependence of ycoefs-> prevents parallelization
              Loop carried backward dependence of ycoefs-> prevents vectorization
              Complex loop carried dependence of Acoefs->,xcoefs-> prevents parallelization
              Generating Tesla code
              30, #pragma acc loop seq
              34, #pragma acc loop vector(128) /* threadIdx.x */
                  Generating implicit reduction(+:sum)
          34, 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.

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

Mot-clé restrict

Une des manières de dire au compilateur que les pointeurs ne sont pas des alias est d'utiliser le mot-clé restrict, 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 __restrict ou __restrict__. Les compilateurs du Portland Group et de NVidia utilisent __restrict. Pour savoir pourquoi il n'existe pas de standard en C++, consultez 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 float * __restrict A; plutôt que float __restrict * A;.

Utilisation du mot-clé restrict

En déclarant un pointeur comme étant restreint, on s'assure qu'uniquement ce pointeur ou une valeur dérivée (comme ptr +1) 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 restrict.

Boucle avec clause independent

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 independent. 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 kernels, nous avons :

#pragma acc kernels
{
#pragma acc loop independent
for (int i=0; i<N; i++)
{
  C[i] = A[i] + B[i];
}
}

Produit matrice-vecteur

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

  double *Acoefs=A.coefs;
  double *xcoefs=x.coefs;
  double *ycoefs=y.coefs;

par le code

  double *__restrict Acoefs=A.coefs;
  double *__restrict xcoefs=x.coefs;
  double *__restrict ycoefs=y.coefs;

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 :

[name@server ~]$ nvc++ -fast -Minfo=accel -ta=tesla:managed main.cpp -o challenge
matvec(const matrix &, const vector &, const vector &):
     23, include "matrix_functions.h"
          27, Generating implicit copyout(ycoefs[:num_rows]) [if not already present]
              Generating implicit copyin(xcoefs[:],row_offsets[:num_rows+1],Acoefs[:],cols[:]) [if not already present]
          30, Loop is parallelizable
              Generating Tesla code
              30, #pragma acc loop gang /* blockIdx.x */
              34, #pragma acc loop vector(128) /* threadIdx.x */
                  Generating implicit reduction(+:sum)
          34, Loop is parallelizable

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 nœud GPU produit ceci :

[name@server ~]$ ./cg.x 
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

Voici le résultat pour la version OpenACC :

[name@server ~]$ ./challenge 
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

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 (nvvp) pour voir ce qui se passe.

1. Démarrez nvvp avec la commande nvvp &, où le symbole & permet de démarrer en arrière-plan.
2. Sélectionnez File -> New Session.
3. Dans le champ "File", cherchez l'exécutable; dans notre exemple, nous utilisons challenge.
4. Cliquez sur "Next" jusqu'à ce que vous puissiez cliquer sur "Finish".

NVIDIA Visual Profiler

NVVP profiler

Browse for the executable you want to profile

One graphical profiler available for OpenACC applications is the NVIDIA Visual Profiler (NVVP). It's a cross-platform analyzing tool for codes written with OpenACC and CUDA C/C++ instructions. Consequently, if the executable is not using the GPU, you will get no result from this profiler.

When X11 is forwarded to an X-Server, or when using a Linux desktop environment (also via JupyterHub with two (2) CPU cores, 5000M of memory and one (1) GPU), it is possible to launch the NVVP from a terminal:

[name@server ~]$ module load cuda/11.7 java/1.8

[name@server ~]$ nvvp

1. After the NVVP startup window, you get prompted for a Workspace directory, which will be used for temporary files. Replace home with scratch in the suggested path. Then click OK.
2. Select File > New Session, or click on the corresponding button in the toolbar.
3. Click on the Browse button at the right of the File path editor.
   1. Change directory if needed.
   2. Select an executable built from codes written with OpenACC and CUDA C/C++ instructions.
4. Below the Arguments editor, select the profiling option Profile current process only.
5. Click Next > to review additional profiling options.
6. Click Finish to start profiling the executable.

This can be done with the following steps:

1. Start nvvp with the command nvvp & (the & sign is to start it in the background)
2. Go in File -> New Session
3. In the "File:" field, search for the executable (named challenge in our example).
4. Click "Next" until you can click "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, Mouvement des données.

La directive parallel loop

Avec la directive kernels, 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 parallel qui peut être combinée à la directive loop ainsi :

#pragma acc parallel loop
for (int i=0; i<N; i++)
{
  C[i] = A[i] + B[i];
}

Comme parallel loop est une directive prescriptive, le compilateur est forcé d'exécuter la boucle en parallèle. Ceci signifie que la clause independent mentionnée plus haut est implicite à l'intérieur d'une zone parallèle.

Pour utiliser cette directive dans notre exemple du produit matrice-vecteur, nous avons besoin des clauses private et reduction pour gérer le flux des données dans la zone parallèle.

• Avec la clause private, une copie de la variable est faite pour chaque itération de la boucle; la valeur de la variable est ainsi indépendante des autres itérations.
• Avec la clause reduction, 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).

Ces clauses ne sont pas nécessaires avec la directive kernels puisque celle-ci fait le travail pour vous.

Reprenons l'exemple de produit matrice-vecteur avec la directive parallel loop;

#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;
  }

La compilation produit le message suivant :

[name@server ~]$ pgc++ -fast -Minfo=accel -ta=tesla:managed main.cpp -o challenge
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

Différences entre parallel loop et kernels

PARALLEL LOOP	KERNELS
l'intégrité du code parallélisé revient au programmeur le programmeur voit la parallélisation que le compilateur pourrait manquer le fonctionnement est identique en OpenMP	la responsabilité d'analyser le code et de garantir son intégrité revient au compilateur une seule directive peut s'appliquer à une grande portion de code le compilateur est libre d'optimiser le code

Les deux approches sont valides et leur performance est comparable.

Exercice : utiliser kernels ou parallel loop

1. Modifiez les fonctions matvec, waxpby et dot. Vous pouvez utiliser soit kernels, soit parallel loop. La solution se trouve dans les répertoires step1. de Github *.
2. Modifiez le Makefile en ajoutant -ta=tesla:managed et -Minfo=accel aux indicateurs pour le compilateur.

Page suivante, Mouvement des données
Retour au début du tutoriel