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.
mbr из 2-х полигонов проверяется, чтобы увидеть, перекрываются ли они, и функция "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

Спасибо за помощь

Answer 1

Когда в предложении данных OpenACC используются агрегированные типы данных, выполняется поверхностная копия типа. Скорее всего, здесь происходит то, что, когда массивы listOfPolygons копируются на устройство, "mbr" будет содержать адреса хоста. Следовательно, программа выдаст ошибку недопустимого адреса при обращении к "mbr".

Учитывая, что в комментарии говорится, что "mbr" всегда будет 4, самое простое, что нужно сделать, - это сделать "mbr" массив фиксированного размера размером 4.

Предполагая, что вы используете компиляторы PGI с устройством NVIDIA, второй метод заключается в использовании унифицированной памяти CUDA путем компиляции "-ta=tesla:managed". Вся динамическая память будет обрабатываться средой выполнения CUDA и позволять получать доступ к адресам хоста на устройстве. Предупреждение о том, что она доступна только для динамических данных, вся ваша программа может использовать только столько памяти, сколько доступно на устройстве, и это может замедлить вашу программу. http://www.pgroup.com/lit/articles/insider/v6n2a4.htm

Третий вариант - выполнить глубокое копирование агрегатного типа на устройство. Я могу опубликовать пример, если вы решите пойти по этому пути. Я также говорю о предмете в рамках презентации, которую я сделал на GTC2015: https://www.youtube.com/watch?v=rWLmZt_u5u4

Answer 2

Вот упрощенный пример. Ключ заключается в том, чтобы использовать неструктурированные области данных в тех же местах, где вы размещаете данные хоста. Сначала выделите массив структур и создайте или скопируйте массив на устройство. Здесь я просто создаю массив, чтобы данные устройства были мусором, но если бы я сделал копирование, то получилось бы поверхностное копирование, и адреса хоста для "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