OpenACC и объектно-ориентированный C++

Question

Я пытаюсь написать объектно-ориентированный код C++, который распараллелен с OpenACC. Мне удалось найти несколько вопросов о стекопереработке и обсуждениях GTC по OpenACC, но я не смог найти реальных примеров объектно-ориентированного кода.

В этом вопросе пример для OpenACCArray было показано, что делает некоторое управление памятью в фоновом режиме (код доступен на http://www.pgroup.com/lit/samples/gtc15_S5233.tar). Однако мне интересно, возможно ли создать класс, который управляет массивами на более высоком уровне. Например

struct Data
{

//    OpenACCArray<float> a;

    OpenACCArray<Vector3<float>> a3;

    Data(size_t len) {
#pragma acc enter data copyin(this)
//        a.resize(len);
        a3.resize(len);
    }
    ~Data() {
#pragma acc exit data delete(this)
    }
    void update_device() {
//        a.update_device();
        a3.update_device();
    }
    void update_host() {
//        a.update_host();
        a3.update_host();
    }
};

int main(int argc, char *argv[])
{
    const size_t len = 32*128;
    Data d(len);

    d.update_device();
 #pragma acc kernels loop independent present(d)
    for (int i=0; i < len; ++i) {
     float val = (float)i/(float)len;

     d.a3[i].x = val;
     d.a3[i].y = i;
     d.a3[i].z = d.a3[i].x / d.a3[i].y;
    }
    d.update_host();
    for (int i=0; i < len/128; ++i) {
       cout << i << ": " << d.a3[i].x << "," << d.a3[i].y << "," << d.a3[i].z << endl;
    }
    cout << endl;
    return 0;
}

Интересно эта программа работает, но как только я раскомментирую OpenACCArray<float> a;, то есть добавление другого члена в эту структуру данных, я получаю ошибки памяти.FATAL ERROR: variable in data clause is partially present on the device,

Так как OpenACCArray struct - это плоская структура, которая обрабатывает косвенные указатели сама по себе, она должна работать, чтобы скопировать ее как член? Или должен быть указатель на структуру, а указатели должны быть жестко связаны с директивами? Тогда я боюсь, что мне придется использовать указатели псевдонимов, предложенные Джеффом Ларкиным в вышеупомянутом вопросе. Я не против выполнить эту работу, но я не могу найти ссылку, как это сделать. Использование директив компилятора keepgpu,keepptx помогает немного понять, что делает компилятор, но я бы предпочел альтернативу обратному проектированию сгенерированного кода ptx.

Любые указатели на полезные справочные проекты или документы высоко ценятся.

Answer 1

В заголовке OpenACCArray1.h удалите две прагмы "#pragma acc enter data create(this)". Происходит то, что конструктор "Data" создает на устройстве объекты "a" и "a3". Следовательно, когда вторая входная область данных встречается в конструкторе OpenACCArray, устройство этого указателя уже существует.

Это работает, когда есть только один элемент данных, так как "a3" и "Data" совместно используют один и тот же адрес для указателя this. Следовательно, когда встречается вторая прагма ввода данных, текущая проверка обнаруживает, что она уже есть на устройстве, поэтому не создает ее снова. Когда добавляется "a", размер "Data" вдвое больше, чем "a", следовательно, настоящая проверка видит, что указатель this уже существует, но имеет размер, отличный от предыдущего. Вот что означает "частично присутствующая" ошибка. Данные есть, но их размер отличается от ожидаемого.

Только родительский класс / структура должны создавать указатель this на устройстве.

Надеюсь, это поможет, Мат