openACC構造体のリストを渡す

Question

2組のポリゴンが重なっているかどうかを調べるCプログラムがあります。私は

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 

}; 

：ユーザ入力2点のポリゴンの集合（データの各セットは数千のポリゴンを持つ）、プログラムは、私はこのような2構造体を有するSET2

のどのポリゴンとSET1重複のどのポリゴン参照します次のコードセグメント：

#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」彼らがある場合、そうでない場合は、0を1に戻し、listOfBooleanする値を入れ

問題
コードはコンパイルできますが実行できません。これは、次のエラーを返します。

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"の戻り値を "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に集計データ型がOpenACCに使用されている場合は

Answer 1

を支援するための

感謝を実行していないとき、プログラムは問題ありませんデータ節では、型の浅いコピーが実行されます。ここで起こりそうなのは、listOfPolygons配列がデバイスにコピーされるときに、 "mbr"にホストアドレスが含まれることです。したがって、 "mbr"がアクセスされるとプログラムは不正なアドレスエラーを出します。コメントは「MBR」は常に4になり、最も簡単な方法は、「MBR」サイズの固定サイズアレイ4

あなたはNVIDIAデバイスとPGIコンパイラを使用していると仮定しますだという考える

2番目の方法は、 "-ta = tesla：managed"をコンパイルしてCUDA Unified Memoryを使用する方法です。すべてのダイナミックメモリはCUDAランタイムによって処理され、デバイス上でホストアドレスにアクセスできるようになります。ダイナミックデータの場合にのみ使用できるという警告がありますが、プログラム全体ではデバイス上で使用可能なメモリを使用でき、プログラムが遅くなる可能性があります。 http://www.pgroup.com/lit/articles/insider/v6n2a4.htm

3番目のオプションは、集約タイプのディープコピーをデバイスに実行することです。あなたがこのルートに行くことを決めたら、私は例を掲示することができます。私はGTC2015で行ったプレゼンテーションの一部としてこのテーマについても話します。https://www.youtube.com/watch?v=rWLmZt_u5u4

Answer 2

ここでは簡単な例を示します。キーは、ホストデータを割り当てる同じ場所で非構造化データ領域を使用することです。最初に構造体の配列を割り当て、配列をデバイスに作成またはコピーします。ここで私はデバイスデータがガベージであるように配列を作成しますが、もし私がcopyinを実行すると、浅いコピーが発生し、 "mbr"のホストアドレスがデバイスにコピーされます。これを修正するには、デバイス上にそれぞれの "mbr"を作成する必要があります。コンパイラは、デバイス "mbr"ポインタを "attach"して、ゴミ/ホストポインタの値を上書きします。 "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