Déplacement de données OpenACC

Question

Je suis très nouveau à OpenACC et je ne comprends pas très bien le mouvement des données et la clause "#pragma acc data".

J'ai un programme écrit en C. Un extrait du code est comme ça:

#pragma acc data create(intersectionSet[0:intersectionsCount][0:4]) // line 122 
#pragma acc kernels // line 123 
for (int i = 0; i<intersectionsCount; i++){ // line 124 
    intersectionSet[i][0] = 9; // line 125 
} 

intersectionsCount a une valeur 210395. Après avoir compilé et exécutez le code ci-dessus par ce qui suit:

pgcc -o rect_openacc -fast -Minfo -acc -ta=nvidia,time rect.c 

J'ai cette sortie:

time(us): 1,475,607 
122: data region reached 1 time 
    31: kernel launched 210395 times 
     grid: [1] block: [128] 
     device time(us): total=1,475,315 max=15 min=7 avg=7 
     elapsed time(us): total=5,451,647 max=24,028 min=24 avg=25 
123: compute region reached 1 time 
    124: kernel launched 1 time 
     grid: [1644] block: [128] 
     device time(us): total=292 max=292 min=292 avg=292 
     elapsed time(us): total=312 max=312 min=312 avg=312 
156: data region reached 1 time 

J'ai quelques questions après avoir lu la sortie:

1. Je ne sais pas pourquoi c'est dit ligne 31, car la ligne 31 n'a pas de pragma acc. Cela signifie-t-il quelque chose que je ne peux pas retracer?
2. Dans la ligne "31: noyau lancé 210395 fois", il a dit qu'il a lancé 210395 fois le noyau. Je ne sais pas s'il est normal que le noyau ait besoin de se lancer tant de fois, car cette partie a pris 5 451 647 (nous) et je pense que c'est un peu long. Je pense que le for-loop est simple et ne devrait pas prendre autant de temps. Est-ce que j'utilise le pragma d'une mauvaise façon?

Mise à jour
J'ai quelques fichiers d'en-tête du programme. Mais ces fichiers n'ont pas de pragma "acc data" ou "acc kernels".

Après avoir compilé le code avec "-Minfo = all", le résultat se présente comme suit:

breakStringToCharArray: 
11, include "stringHelper.h" 
     50, Loop not vectorized/parallelized: contains call 
countChar: 
11, include "stringHelper.h" 
     74, Loop not vectorized/parallelized: not countable 
extractCharToIntRequiredInt: 
11, include "stringHelper.h" 
     93, Loop not vectorized/parallelized: contains call 
extractArray: 
12, include "fileHelper.h" 
     49, Loop not vectorized/parallelized: contains call 
isRectOverlap: 
13, include "shapeHelper.h" 
     23, Generating acc routine vector 
      Generating Tesla code 
getRectIntersection: 
13, include "shapeHelper.h" 
     45, Generating acc routine vector 
      Generating Tesla code 
getRectIntersectionInGPU: 
13, include "shapeHelper.h" 
     69, Generating acc routine vector 
      Generating Tesla code 
max: 
13, include "shapeHelper.h" 
     98, Generating acc routine vector 
      Generating Tesla code 
min: 
13, include "shapeHelper.h" 
    118, Generating acc routine vector 
      Generating Tesla code 
main: 
64, Loop not vectorized/parallelized: contains call 
108, Loop not vectorized/parallelized: contains call 
122, Generating create(intersectionSet[:intersectionsCount][:4]) 
124, Loop is parallelizable 
    Accelerator kernel generated 
    Generating Tesla code 
124, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */ 

Je crée intersectionSet ainsi:

intersectionSet = (int **)malloc(sizeof(int **) * intersectionsCount); 
for (i = 0; i<intersectionsCount; i++){ 
    intersectionSet[i] = (int *)malloc(sizeof(int *) * 4); 
} 

Answer 1

Qu'est-ce qui se passe est que, puisque vous avez pointeur vers un tableau de pointeurs, "**", (du moins je suppose que c'est intersectionSet) le compilateur doit d'abord allouer le pointeur au pointeur sur le périphérique, puis faire une boucle sur chaque élément pour allouer les tableaux de périphériques individuels. Enfin, il doit ensuite lancer un noyau pour définir la valeur du pointeur sur le périphérique. Voici un pseudo-code pour aider à illustrer.

devPtrPtr = deviceMalloc(numElements*pointer size); 
for (i=0; i < numElements; ++i) { 
    devPtr = deviceMalloc(elementSize * dataTypeSize); 
    call deviceKernelToSetPointer<<<1,128>>(devPtrPtr[i],devPtr); 
} 

Pour aider votre code, je commute les dimensions faisant la longueur de la colonne 4 et la longueur de la ligne « intersectionsCount ». Cela facilitera également l'accès aux données sur l'appareil puisque la boucle "vector" doit correspondre à la dimension stride-1 (contiguë) afin d'éviter les divergences de mémoire.

Hope this helps,

Mat