OpenACC против C++: FATAL ERROR: переменная частично присутствует на устройстве

Здесь происходит то, что адрес хоста «a» совпадает с начальным адресом «DK». Следовательно, когда компилятор ищет адрес хоста в текущей таблице, которую он использует для сопоставления адреса хоста переменной с адресом устройства, он видит, что размер другой. «a» — размер 8, а «DK» — размер 24.

Я покажу исправление ниже, но давайте вернемся назад и поймем, что здесь происходит. Когда «DK» создается на хосте, он сначала создает хранилище для каждого из своих членов данных, а затем вызывает конструктор класса каждого члена данных. Затем он выполняет свой собственный конструктор. Следовательно, для каждого члена данных ваш код создаст указатель класса this на устройстве, а затем выделит массив «arr» на устройстве. После этого на устройстве создается «DK» с пространством для каждого члена данных. Однако, поскольку копия устройства "DK" создается после элементов данных, компилятор не может автоматически связать их.

Ниже я разместил два возможных исправления.

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

Во-вторых, вы можете использовать класс DataKeeper для управления данными класса Data. Однако для этого потребуется, чтобы данные Data были общедоступными.

Обратите внимание, что я написал главу 5 «Расширенное управление данными» книги «Параллельное программирование с OpenACC» и включил в нее раздел об управлении данными классов C++. Вы можете найти мой пример кода по адресу: https://github.com/rmfarber/ParallelProgrammingWithOpenACC/tree/master/Chapter05 В частности, посмотрите, как я создал общий класс-контейнер "accList".

Исправление №1:

#define DATASIZE 16
#include <iostream>
#ifdef _OPENACC
#include <openacc.h>
#endif

class Data {
  float *arr;

public:
  Data() {
    arr = new float[DATASIZE];
#pragma acc enter data copyin(this)
#pragma acc enter data create(arr[:DATASIZE])
  }

  ~Data() {
#pragma acc exit data delete(arr)
#pragma acc exit data delete(this)
    delete [] arr;
  }

  float &get(int i) { return arr[i]; }
  void updatehost() {
   #pragma acc update host(arr[0:DATASIZE])
  }

};

class DataKeeper {
  Data *a, *b, *c;

public:
  DataKeeper() {
#pragma acc enter data copyin(this)
  a = new Data;
  b = new Data;
  c = new Data;
  }

  ~DataKeeper() {
#pragma acc exit data delete(this)
  delete a;
  delete b;
  delete c;
  }

  void init() {
#pragma acc parallel loop present(a,b,c)
    for (int i = 0; i < DATASIZE; ++i) {
      a->get(i) = i;
    }
    a->updatehost();
    std::cout << "a.arr[0]=" << a->get(0) << std::endl;
    std::cout << "a.arr[end]=" << a->get(DATASIZE-1) << std::endl;
  }
};

int main() {
  DataKeeper DK;
  DK.init();
}

Исправление № 2

#define DATASIZE 16
#include <iostream>
#ifdef _OPENACC
#include <openacc.h>
#endif

class Data {
public:
  float *arr;

  Data() {
    arr = new float[DATASIZE];
  }

  ~Data() {
    delete [] arr;
  }

  float &get(int i) { return arr[i]; }
};

class DataKeeper {
  Data a, b, c;

public:
  DataKeeper() {
#pragma acc enter data copyin(this)
#pragma acc enter data create(a.arr[0:DATASIZE])
#pragma acc enter data create(b.arr[0:DATASIZE])
#pragma acc enter data create(c.arr[0:DATASIZE])
  }

  ~DataKeeper() {
#pragma acc exit data delete(this)
#pragma acc exit data delete(a.arr)
#pragma acc exit data delete(b.arr)
#pragma acc exit data delete(c.arr)
  }

  void init() {
#pragma acc parallel loop present(a,b,c)
    for (int i = 0; i < DATASIZE; ++i) {
      a.get(i) = i;
    }
#pragma acc update host(a.arr[0:DATASIZE])
    std::cout << "a.arr[0]=" << a.arr[0] << std::endl;
    std::cout << "a.arr[end]=" << a.arr[DATASIZE-1] << std::endl;
  }
};

int main() {
  DataKeeper DK;
  DK.init();
}