openACC 传递结构列表

Question

我有一个 C 程序来查找 2 组多边形是否重叠。用户输入2组多边形（每组数据有数千个多边形），程序查看set1中的哪个多边形与set2中的哪个多边形重叠

我有 2 个这样的结构：

struct gpc_vertex  /* Polygon vertex */
{
    double          x;
    double          y;
};

struct gpc_vertex_list  /* Polygon contour */
{
    int pid;    // polygon id
    int             num_vertices;
    double *mbr;   // minimum bounding rectangle of the polygon, so always 4 elements

};

我有以下代码段：

#pragma acc kernels copy(listOfPolygons1[0:polygonCount1], listOfPolygons2[0:polygonCount2], listOfBoolean[0:dump])
for (i=0; i<polygonCount1; i++){
    polygon1 = listOfPolygons1[i];

    for (j=0; j<polygonCount2; j++){

        polygon2 = listOfPolygons2[j];
        idx = polygonCount2 * i + j;

        listOfBoolean[idx] = isRectOverlap(polygon1.mbr, polygon2.mbr);  // line 115

    }
}

listOfPolygons1 和 listOfPolygons2（顾名思义）是 gpc_vertex_list 的数组。 listOfBoolean 是一个 int 数组。
检查 2 个多边形的 mbr 以查看它们是否重叠，函数 "isRectOverlap" return 如果重叠则为 1，如果不重叠则为 0 并将值放入 listOfBoolean

问题
代码可以编译但不能运行。它 return 出现以下错误：

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

我的观察
该程序可以通过将第 115 行更改为以下内容来编译和运行：

isRectOverlap(polygon1.mbr, polygon2.mbr); // without assigning value to listOfBoolean

或者这个：

listOfBoolean[idx] = 5; // assigning an arbitrary value

(虽然结果不对，但至少，可以运行)

问题
如果值没有从 "isRectOverlap" 传递到 "listOfBoolean"
，"isRectOverlap" 和 "listOfBoolean" 似乎都不会产生问题有谁知道如果我将 "isRectOverlap" 中的 return 值分配给 "listOfBoolean" 为什么它不能运行？

isRectOverlap函数是这样的：

int isRectOverlap(double *shape1, double *shape2){

    if (shape1[0] > shape2[2] || shape2[0] > shape1[2]){
        return 0;
    }

    if (shape1[1] < shape2[3] || shape2[1] < shape1[3]){
        return 0;
    }

    return 1;

}

不运行ning在OpenACC

下程序没有问题

感谢您的帮助

Answer 1

当在 OpenACC 数据子句中使用聚合数据类型时，将执行该类型的浅表复制。此处最有可能发生的情况是，当将 listOfPolygons 数组复制到设备时，"mbr" 将包含主机地址。因此，程序在访问"mbr"时会报非法地址错误。

鉴于评论说 "mbr" 将始终为 4，最简单的做法是使 "mbr" 成为大小为 4 的固定大小数组。

假设您在 NVIDIA 设备上使用 PGI 编译器，第二种方法是通过编译“-ta=tesla:managed”来使用 CUDA 统一内存。所有动态内存都将由 CUDA 运行时处理，并允许在设备上访问主机地址。需要注意的是它只适用于动态数据，你的整个程序只能使用设备上可用的内存，这可能会减慢你的程序。 http://www.pgroup.com/lit/articles/insider/v6n2a4.htm

第三种选择是对设备执行聚合类型的深层复制。如果你决定走这条路，我可以 post 举个例子。我还在 GTC2015 做的演讲中谈到了这个主题：https://www.youtube.com/watch?v=rWLmZt_u5u4

Answer 2

这是一个简化的例子。关键是在分配主机数据的相同位置使用非结构化数据区域。首先分配结构数组，然后创建或复制数组到设备。在这里，我只是创建了一个数组，所以设备数据是垃圾，但如果我做了一个复制，那么就会发生一个浅拷贝，"mbr" 的主机地址将被复制到设备。要解决此问题，您需要在设备上创建每个 "mbr"。然后，编译器将分配 "attach" 设备 "mbr" 指针，从而覆盖 garbage/host 指针值。一旦 "mbr" 具有有效的设备指针，就可以在设备上引用它们。

% cat example_struct.c
#include <stdlib.h>
#include <stdio.h>
#ifndef N
#define N 1024
#endif

typedef struct gpc_vertex_list
{
    int pid;    // polygon id
    int num_vertices;
    double *mbr;   // minimum bounding rectangle of the polygon, so always 4 elements

} gpc_vertex_list;

gpc_vertex_list * allocData(size_t size);
int deleteData(gpc_vertex_list * A, size_t size);
int initData(gpc_vertex_list *Ai, size_t size);

#pragma acc routine seq
int isRectOverlap(double * mbr) {
    int result;
    result = mbr[0];
    result += mbr[1];
    result += mbr[2];
    result += mbr[3];
    return result;
}

int main() {
    gpc_vertex_list *A;
    gpc_vertex_list B;
    size_t size, i;
    int * listOfBoolean;
    size = N;
    A=allocData(size);
    initData(A,size);
    listOfBoolean = (int*) malloc(sizeof(int)*size);

#pragma acc parallel loop present(A) copyout(listOfBoolean[0:size])  private(B)
    for (i=0; i<size; i++){
       B = A[i];
       listOfBoolean[i] = isRectOverlap(B.mbr);
    }

    printf("result: %d %d %d\n",listOfBoolean[0], listOfBoolean[size/2], listOfBoolean[size-1]);
    free(listOfBoolean);
    deleteData(A, size);
    exit(0);
}

gpc_vertex_list * allocData(size_t size) {
    gpc_vertex_list * tmp;
    tmp = (gpc_vertex_list *) malloc(size*sizeof(gpc_vertex_list));
/* Create the array on device.  */
#pragma acc enter data create(tmp[0:size])
    for (int i=0; i< size; ++i) {
       tmp[i].mbr = (double*) malloc(sizeof(double)*4);
/* create the member array on the device */
#pragma acc enter data create(tmp[i].mbr[0:4])
    }
    return tmp;
}

int deleteData(gpc_vertex_list * A, size_t size) {
/* Delete the host copy. */
    for (int i=0; i< size; ++i) {
#pragma acc exit data delete(A[i].mbr)
        free(A[i].mbr);
    }
#pragma acc exit data delete(A)
    free(A);
}

int initData(gpc_vertex_list *A ,size_t size) {
    size_t i;
    for (int i=0; i< size; ++i) {
       A[i].pid = i;
       A[i].num_vertices = 4;
       for (int j=0; j<4;++j) {
           A[i].mbr[j]=(i*4)+j;
       }
       #pragma acc update device(A[i].pid,A[i].num_vertices,A[i].mbr[0:4])
    }
}
% pgcc example_struct.c -acc -Minfo=accel
isRectOverlap:
     20, Generating acc routine seq
main:
     39, Generating copyout(listOfBoolean[:size])
         Generating present(A[:])
         Accelerator kernel generated
         Generating Tesla code
         40, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     39, Local memory used for B
allocData:
     55, Generating enter data create(tmp[:size])
     59, Generating enter data create(tmp->mbr[:4])
deleteData:
     67, Generating exit data delete(A->mbr[:1])
     70, Generating exit data delete(A[:1])
initData:
     83, Generating update device(A->mbr[:4],A->pid,A->num_vertices)
% a.out
result: 6 8198 16374

openACC 传递结构列表

openACC passing a list of struct

c

struct

openacc

pgi

pgcc