Utilisation d’OpenACC pour paralléliser des boucles nestedes

Je suis très nouveau pour openacc et n’ai que des connaissances de haut niveau, toute aide et explication de ce que je fais de mal serait appréciée.

J’essaie d’accélérer (paralléliser) une boucle nestede pas si simple qui met à jour un tableau aplati (3D à 1D) à l’aide de directives openacc. J’ai posté un exemple de code simplifié ci-dessous qui, une fois compilé à l’aide de

pgcc -acc -Minfo=accel test.c

donne l’erreur suivante:

call to cuStreamSynchronize returned error 700: Illegal address during kernel execution

Code:

 #include  #include  #define min(a,b) (a > b) ? b : a #define max(a,b) (a distance[0:NX*NY*NZ]) for(int level = startLevel; level <= endLevel; level++){ int ks = max(1, level-(y + z)); int ke = min(x, level-2); int js = max(1, level-(x + z)); int je = min(y, level-2); #pragma acc region { #pragma acc loop independent for(int k = ks; k <= ke; k++){ #pragma acc loop independent for(int j = js; j  0 && i distance[index] = solve(p, index); } } } } } } void create_phi(Phi *p){ p->dx = 1; p->dy = 1; p->dz = 1; p->distance = (double *) malloc(sizeof(double) * NX * NY * NZ); for(int i = 0; i < NZ; i++){ for(int j = 0; j < NY; j++){ for(int k = 0; k distance[index] = (i*j*k == 0) ? 0 : 1; } } } } int main() { printf("start \n"); Phi *p = (Phi *) malloc(sizeof(Phi)); create_phi(p); printf("calling fast sweep \n"); fast_sweep(p); printf(" print the results \n"); for(int i = 0; i < NZ; i++){ for(int j = 0; j < NY; j++){ for(int k = 0; k distance[index]); } printf("\n"); } printf("\n"); } return 0; } 

Au lieu d’utiliser les directives region et loop , utilisez les commandes

 #pragma acc kernels 

produit l’erreur suivante:

 solve: 19, Generating acc routine seq fast_sweep: 34, Generating copy(p->distance[:1000]) 42, Generating copy(p[:1]) 45, Loop carried dependence due to exposed use of p[:1] prevents parallelization Accelerator scalar kernel generated 47, Loop carried dependence due to exposed use of p[:i1+1] prevents parallelization 

Je lance ce code sur

 GNU/Linux CentOS release 6.7 (Final) GeForce GTX Titan pgcc 15.7-0 64-bit target on x86-64 Linux -tp sandybridge 

    L’erreur provient du kernel de calcul sur le GPU qui déréférence un pointeur de CPU. C’est un problème assez courant que le comité OpenACC s’emploie à résoudre. Des structures de données dynamics telles que celles-ci peuvent poser de nombreux problèmes, nous voulons donc les résoudre. Voici deux solutions possibles pour vous.

    1) Utilisez l’option «Mémoire gérée» via l’option «Package d’évaluation de la mémoire unifiée» de PGI lors de l’installation du compilateur. Il s’agit d’une fonctionnalité bêta , mais toutes vos données seront stockées dans un type de mémoire spécial, visible à la fois pour le processeur et le processeur graphique. Il y a beaucoup de mises en garde que vous devriez lire dans la documentation, notamment que vous êtes limité à la quantité de mémoire disponible sur le GPU et que vous ne pouvez pas accéder à la mémoire du processeur lorsqu’elle est utilisée sur le GPU. une solution de contournement possible. En supposant que vous ayez activé cette option lors de l’installation, ajoutez simplement -ta=tesla:managed aux indicateurs de votre compilateur pour l’activer. J’ai essayé cela avec votre code et cela a fonctionné.

    2) Ajoutez un pointeur sur votre code pour que vous n’accédiez pas à la distance via p , mais que vous y accédiez directement, comme suit:

     double *distance = p->distance; #pragma acc data copy(p[0:1],distance[0:NX*NY*NZ]) for(int level = startLevel; level <= endLevel; level++){ int ks = max(1, level-(y + z)); int ke = min(x, level-2); int js = max(1, level-(x + z)); int je = min(y, level-2); #pragma acc parallel { #pragma acc loop independent for(int k = ks; k <= ke; k++){ #pragma acc loop independent for(int j = js; j <= je; j++){ int i = level - (k + j); if(i > 0 && i <= z){ int index = i * NX * NY + j * NX + k; distance[index] = solve(p, index); } } } } 

    Je sais que cela peut être pénible lorsque de nombreux tableaux de données sont utilisés, mais c'est une solution de contournement que j'ai utilisée avec succès dans de nombreux codes. C’est regrettable que cela soit nécessaire, c’est pourquoi nous souhaitons proposer une meilleure solution dans une future version d’OpenACC.

    J'espère que ça aide! Si je peux trouver une solution qui ne nécessite pas de pointeur supplémentaire, je mettrai à jour cette réponse.

    Jeff a raison de dire que le comité OpenACC travaille toujours sur la manière de normaliser la prise en charge des types de données agrégés avec des membres de données dynamics. Cependant, avec la version 14.9 ou ultérieure de PGI, nous avons ajouté une meilleure prise en charge des structures ainsi que des classes C ++, ce qui vous permet de simplifier le code en ajoutant simplement create(p[0:1]) . Ce qui va arriver, c’est que le compilateur créera une copie de périphérique de p avec de la mémoire allouée uniquement aux membres de données. Ensuite, lorsque vous faites la copie de p->distance , la mémoire est allouée pour “distance”, puis reliez-la à p . (C’est-à-dire que le temps d’exécution renseignera le pointeur de périphérique dans la structure)

    Il y a des mises en garde. Premièrement, ce comportement n’a pas été normalisé. Par conséquent, d’autres compilateurs tels que Cray, Pathscale, GCC et d’autres peuvent avoir un comportement différent. Deuxièmement, l’ordre est important. p doit être créé avant que la distance puisse être fixée. Troisièmement, les structures de données plus complexes deviennent très difficiles à gérer. Comme Jeff le suggère, l’utilisation de la mémoire unifiée CUDA est une bonne alternative pour gérer des structures de données complexes.

    Si vous êtes intéressé, une grande partie de ma présentation GTC2015 traite de ce sujet ( lien ). L’exposé porte sur la gestion des données de classe C ++, mais s’applique également aux structures C.

    J’espère que ça aide, Mat

     % cat test1.c #include  #include  #define min(a,b) (a > b) ? b : a #define max(a,b) (a < b) ? b : a #define NX 10 #define NY 10 #define NZ 10 struct phiType { double dx, dy, dz; double * distance; }; typedef struct phiType Phi; #pragma acc routine seq double solve(Phi *p, int index) { // for simplicity just returning a value return 2; } void fast_sweep(Phi *p) { // removing boundaries int x = NX - 2; int y = NY - 2; int z = NZ - 2; int startLevel = 3; int endLevel = x + y + z; #pragma acc data create(p[0:1]) copy(p->distance[0:NX*NY*NZ]) for(int level = startLevel; level <= endLevel; level++){ int ks = max(1, level-(y + z)); int ke = min(x, level-2); int js = max(1, level-(x + z)); int je = min(y, level-2); #pragma acc region { #pragma acc loop independent for(int k = ks; k <= ke; k++){ #pragma acc loop independent for(int j = js; j <= je; j++){ int i = level - (k + j); if(i > 0 && i <= z){ int index = i * NX * NY + j * NX + k; p->distance[index] = solve(p, index); } } } } } } void create_phi(Phi *p){ p->dx = 1; p->dy = 1; p->dz = 1; p->distance = (double *) malloc(sizeof(double) * NX * NY * NZ); for(int i = 0; i < NZ; i++){ for(int j = 0; j < NY; j++){ for(int k = 0; k < NX; k++){ int index = i * NX * NY + j * NX + k; p->distance[index] = (i*j*k == 0) ? 0 : 1; } } } } int main() { printf("start \n"); Phi *p = (Phi *) malloc(sizeof(Phi)); create_phi(p); printf("calling fast sweep \n"); fast_sweep(p); printf(" print the results \n"); for(int i = 0; i < NZ; i++){ for(int j = 0; j < NY; j++){ for(int k = 0; k < NX; k++){ int index = i * NX * NY + j * NX + k; printf("%f ", p->distance[index]); } printf("\n"); } printf("\n"); } return 0; } % pgcc -acc -ta=tesla:cc35 -Minfo=accel test1.c -V15.7 ; a.out solve: 19, Generating acc routine seq fast_sweep: 34, Generating create(p[:1]) Generating copy(p->distance[:1000]) 45, Loop is parallelizable 47, Loop is parallelizable Accelerator kernel generated Generating Tesla code 45, #pragma acc loop gang /* blockIdx.y */ 47, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */ start calling fast sweep print the results 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 2.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 0.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000 1.000000