Как инициализировать объект в месте расположения указателя с помощью cudaMallocManaged? (C ++)

Question

Ниже показана упрощенная версия кода, с которым я работаю, когда я не использую cudaMallocManaged и просто инициализирую все, используя ключевое слово new, все работает просто отлично. Я сузил проблему до способа инициализации элементов массива в for l oop, но я не могу понять, что здесь происходит не так. При печати, сравнивая указатели, все выглядит так, как должно, но при этом возникает ошибка сегментации.

#include <iostream>

class B
{
    public:
        __device__ __host__ virtual void test() = 0;
};

class A: public B
{
    public:
        __device__ __host__ A(int x) {number = x;};
        __device__ __host__ void test() {printf("test called!\n");}

        int number;
};

int main(int argc, char const *argv[])
{
    // Size of array.
    static const int count = 2;

    // Create array of pointers to A objects in memmory.
    B** list; // = new B*[count];
    cudaMallocManaged(&list, count*sizeof(B*));

    // Create objects for in array.
    for (int i = 0; i < count; i++)
    {
        A* tempPointer;
        cudaMallocManaged(&tempPointer, sizeof(A));
        *tempPointer = A(500);
        list[i] = tempPointer;
    }

    // Gives a segmentation fault.
    for (int i = 0; i < count; i++)
        list[i]->test();

    // Free memmory.
    for (int i = 0; i < count; i++)
        cudaFree(list[count]);
    cudaFree(list);
}

Использование этого для l oop вместо этого приведет к рабочему коду, но мне действительно нужно использовать cudaMallocManaged, так что это не вариант:

for (int i = 0; i < count; i++)
{
    A* tempPointer = new A(500);
    list[i] = tempPointer;
}

Answer 1

Проблема здесь заключается в том, что способ инициализации объекта класса, содержащего виртуальные методы и, следовательно, таблицу указателей виртуальных функций:

class B
{
    public:
        __device__ __host__ virtual void test() = 0;
};

class A: public B
{
    public:
        __device__ __host__ A(int x) {number = x;};
        __device__ __host__ void test() {printf("test called!\n");}

        int number;
};

, заключается не в копировании объекта:

   *tempPointer = A(500);

Этот метод не инициализирует таблицу указателей виртуальных функций в объекте.

Вместо этого, в этом конкретном случае, я бы рекомендовал использовать новое размещение:

$ cat t1674.cu
#include <iostream>
#include <stdio.h>

class B
{
    public:
        __device__ __host__ virtual void test() = 0;
};

class A: public B
{
    public:
        __device__ __host__ A(int x) {number = x;};
        __device__ __host__ void test() {printf("test called!\n");}

        int number;
};

int main(int argc, char const *argv[])
{
    // Size of array.
    static const int count = 2;

    // Create array of pointers to A objects in memmory.
    B** list; // = new B*[count];
    cudaMallocManaged(&list, count*sizeof(B*));

    // Create objects for in array.
    for (int i = 0; i < count; i++)
    {
        A* tempPointer;
        cudaMallocManaged(&tempPointer, sizeof(A));
//        *tempPointer = A(500);
        list[i] = new(tempPointer) A(500);
    }

    // Gives a segmentation fault.
    for (int i = 0; i < count; i++)
        list[i]->test();

    // Free memmory.
    for (int i = 0; i < count; i++)
        cudaFree(list[count]);
    cudaFree(list);
}
$ nvcc -o t1674 t1674.cu
$ cuda-memcheck ./t1674
========= CUDA-MEMCHECK
test called!
test called!
========= ERROR SUMMARY: 0 errors
$

Примечание Выше я также исправил еще одну ошибку в коде, в частности, что вы пытаетесь освободить указатель list[0] более одного раза, что, очевидно, не может быть правильным. Я изменил его на list[count], который, как я полагаю, был вашим намерением.

Сказав все это, я подозреваю, что вскоре вы столкнетесь с проблемой такого подхода. CUDA имеет ограничения вокруг объектов с таблицами указателей виртуальных функций. В частности, объект должен быть создан в домене, который будет использоваться. Если вы собираетесь использовать его только на хосте, инициализируйте объект на хосте. Если вы собираетесь использовать его только на устройстве, инициализируйте объект на устройстве. Объекты (с таблицами указателей виртуальных функций), инициализированные в одном домене, нельзя безопасно использовать в другом .