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 行没有 acc pragma。这是否意味着我无法追踪的东西？
2. 在“31：内核启动210395次”行中，它表示它启动了210395次内核。我不知道内核需要启动这么多次是否正常，因为这部分花了5,451,647（us），我认为它有点长。我认为for循环很简单，不应该花太多时间。我是否以错误的方式使用了编译指示？

更新
我确实有几个程序的头文件。但这些文件没有“acc data”或“acc kernels”编译指示。

用“-Minfo=all”编译代码后，结果如下：

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”。这也将有助于设备上的数据访问，因为“向量”循环应对应于 stride-1（连续）维度以避免内存分歧。

希望这会有所帮助，

垫子