Движение данных OpenACC

Question

Я очень новичок в OpenACC, и я не совсем понимаю движение данных и предложение «#pragma acc data».

У меня есть программа, написанная на C. Выписка из кода так:

#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 имеет значение 210395. После компиляции и запуска выше код следующим:

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

у меня есть этот выход:

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 

у меня есть несколько вопросов после прочтения вывода:

1. Я не знаю, почему он сказал строку 31, так как линия 31 не имеет соответствующей прагмы. Означает ли это то, что я не могу проследить?
2. В строке «31: ядро ​​запущено 210395 раз», он сказал, что запустил 210395 раз ядро. Я не знаю, нормально ли, что ядро ​​нужно запускать столько раз, потому что эта часть заняла 5 451 647 (нас), и я думаю, что это немного длиннее. Я думаю, что for-loop прост и не должен занимать столько времени. Я использую прагму неправильно?

Update
У меня есть несколько файлов заголовка для программы. Но в этих файлах нет прагмы «acc data» или «acc kernels».

После компиляции кода с «-Minfo = все», результат выглядит следующим образом:

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 */ 

Я создаю intersectionSet таким образом:

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

Answer 1

Что происходит, что поскольку у вас есть указатель на массив указателей, «**» (по крайней мере, я предполагаю, что это что такое intersectionSet), компилятор должен сначала выделить указатель на указатель на устройстве, а затем перебрать каждый элемент для выделения отдельных массивов устройств. Наконец, тогда необходимо запустить ядро ​​для установки значения указателя на устройстве. Вот несколько Псевдокодов, которые помогут проиллюстрировать.

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

Чтобы помочь вашему коду, я бы переключил размеры, составляющие длину столбца 4 и длину строки «intersectionsCount». Это также поможет доступу к данным на устройстве, поскольку цикл «вектор» должен соответствовать размеру шага 1 (смежный), чтобы избежать расхождения в памяти.

Надеется, что это помогает,

Мат