ошибка openacc при присвоении значений динамически выделяемому массиву элементов структуры структуры, на которую ссылается указатель

Question

Я пытаюсь обдумать объединение openacc с указателями на структуры, содержащие динамически распределяемые элементы. Код ниже не работает с

Сбой в потоке:1 вызов cuStreamSynchronize возвратил ошибку 700: недопустимый адрес во время выполнения ядра

при компиляции с использованием nvc («64-разрядная цель nvc 20.9-0 LLVM на x86-64 Linux -tp haswell»). Насколько я могу судить, я следую подходу, предложенному, например, в руководстве по началу работы с OpenACC. Но почему-то предположительно указатели не прилипают (?) к устройству. Кто-нибудь знает, что здесь не так?

      #include <stdlib.h>
#include <stdio.h>

typedef struct grid
{
  int N;
  double *X;
} grid;

void allocate(grid* g, int N)
{
  g->N = N;
  g->X = (double*) malloc(sizeof(double) * g->N);

  #pragma acc enter data create(g[0:1])
  #pragma acc enter data create(g->X[0:N])
}

void release(grid* g)
{
  #pragma acc exit data delete(g->X[0:g->N])
  #pragma acc exit data delete(g[0:1])

  free(g->X);
}

void fill(grid * g)
{
  int i;

  #pragma acc parallel loop
  for (i = 0; i < g->N; i++)
  {
    g->X[i] = 42; // the cuprit, commenting this removes the error too
  }
}

int main()
{
  grid g;

  allocate(&g, 10);

  fill(&g);

  release(&g);

  return 0;
}```

Answer 1

Из сообщений обратной связи компилятора вы увидите что-то вроде:

           fill:
          32, Accelerator restriction: size of the GPU copy of g is unknown
              Generating Tesla code
              32, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
          32, Generating implicit copyin(g) [if not already present]
          37, Generating update self(g->X[:g->N])

Проблема в том, что компилятор не может неявно копировать агрегатные типы с элементами динамических данных, поэтому вам нужно добавить «представление (g)», чтобы указать, что g уже является устройством.

Кроме того, вы захотите скопировать, чтобы получить значение N на устройстве, и вам не нужно включать форму массива в директиву удаления выходных данных. Например:

      % cat test.c
#include <stdlib.h>
#include <stdio.h>

typedef struct grid
{
  int N;
  double *X;
} grid;

void allocate(grid* g, int N)
{
  g->N = N;
  g->X = (double*) malloc(sizeof(double) * g->N);

  #pragma acc enter data copyin(g[0:1])
  #pragma acc enter data create(g->X[0:N])
}

void release(grid* g)
{
  #pragma acc exit data delete(g->X)
  #pragma acc exit data delete(g)

  free(g->X);
}

void fill(grid * g)
{
  int i;

  #pragma acc parallel loop present(g)
  for (i = 0; i < g->N; i++)
  {
    g->X[i] = 42; // the cuprit, commenting this removes the error too
  }
  #pragma acc update self(g->X[:g->N])
  for (i = 0; i < 4; i++)
  {
    printf("%d : %f \n",i,g->X[i]);
  }
}

int main()
{
  grid g;

  allocate(&g, 10);

  fill(&g);

  release(&g);

  return 0;
}

% nvc -acc test.c -Minfo=accel -V20.9 ; a.out
allocate:
     17, Generating enter data copyin(g[:1])
         Generating enter data create(g->X[:N])
release:
     24, Generating exit data delete(g[:1],g->X[:1])
fill:
     32, Generating present(g[:1])
         Generating Tesla code
         32, #pragma acc loop gang, vector(128) /* blockIdx.x threadIdx.x */
     37, Generating update self(g->X[:g->N])
0 : 42.000000
1 : 42.000000
2 : 42.000000
3 : 42.000000