38,760
edits
OpenACC Tutorial - Adding directives/fr: Difference between revisions
OpenACC Tutorial - Adding directives/fr (view source)
Revision as of 22:06, 23 December 2022
, 1 year agoUpdating 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 : | Voici un exemple de cette directive : | ||
| 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]; | |||
} | } | ||
} | } | ||
</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 : | 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 : | ||
<syntaxhighlight lang="cpp" line> | </div> | ||
for(int i=0;i<num_rows;i++) { | <syntaxhighlight lang="cpp" line start="29"> | ||
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> | </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 ==== | ||
<div class="mw-translate-fuzzy"> | |||
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" | ||
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( | Generating implicit copyin(xcoefs[:]) [if not already present] | ||
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 | Generating NVIDIA GPU code | ||
31, #pragma acc loop seq | |||
35, #pragma acc loop vector(128) /* threadIdx.x */ | |||
Generating implicit reduction(+:sum) | Generating implicit reduction(+:sum) | ||
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 == | ||