openACC передает список структур

Когда агрегатные типы данных используются в предложении данных 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

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