1

我有一个 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”返回 1，如果不是，则返回 0，并将值放入 listOfBoolean

问题
代码可以编译但不能运行。它返回以下错误：

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

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

isRectOverlap(polygon1.mbr, polygon2.mbr); // 不给 listOfBoolean 赋值

或这个：

listOfBoolean[idx] = 5; // 分配一个任意值

（虽然结果是错误的，但至少，它可以运行）

问题
如果值没有从“isRectOverlap”传递到“listOfBoolean”，“isRectOverlap”和“listOfBoolean”似乎都不会产生问题
有谁知道如果我将“isRectOverlap”的返回值分配给“ 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;

}

程序不在OpenACC中运行时没有问题

感谢您的帮助

4

2 回答 2

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

第三种选择是对设备执行聚合类型的深拷贝。如果您决定走这条路，我可以发布一个示例。作为我在 GTC2015 上的演讲的一部分，我还谈到了这个主题：https ://www.youtube.com/watch?v=rWLmZt_u5u4

于 2016-08-05T22:03:08.207 回答
1

这是一个简化的例子。关键是在分配主机数据的相同位置使用非结构化数据区域。首先分配结构数组并创建或复制数组到设备。在这里，我只是创建了数组，因此设备数据是垃圾，但是如果我进行了复制，则会发生浅拷贝，并且“mbr”的主机地址将被复制到设备。要解决此问题，您需要在设备上创建每个“mbr”。然后编译器将分配“附加”设备“mbr”指针，从而覆盖垃圾/主机指针值。一旦“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
于 2016-08-08T16:38:24.770 回答