Могу ли я использовать класс, который реализует чисто виртуальные функции внутри ядра CUDA?

Я борюсь с проблемой, которая кажется немного неясной.

Я работаю над фреймворком, в котором пользователь может предоставить реализацию абстрактного базового класса, который после нескольких шагов магии и генерации кода будет использоваться внутри ядра CUDA.

я знаю это

"Не разрешается передавать в качестве аргумента глобальной функции объект класса с виртуальными функциями."

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

class VirtualBase {
public:
    __device__ virtual int getResult() const = 0;
    __device__ virtual ~VirtualBase();
};

class Implementation : public VirtualBase {
public:
    __device__ Implementation(){};
    __device__ int getResult() const { return 42; };
    __device__ ~Implementation() {};
};

__global__ void kernel() {
    Implementation impl;
    int res = impl.getResult();
}

int main(void) {
    kernel<<<1, 1>>>();
    return 0;
}

Код скомпилирован с автоматически сгенерированным make-файлом Nsights.

/Developer/NVIDIA/CUDA-7.5/bin/nvcc -G -g -O0 -std=c++11 -gencode arch=compute_30,code=sm_30  -odir "src" -M -o "src/main.d" "../src/main.cu"
/Developer/NVIDIA/CUDA-7.5/bin/nvcc -G -g -O0 -std=c++11 --compile --relocatable-device-code=false -gencode arch=compute_30,code=compute_30 -gencode arch=compute_30,code=sm_30  -x cu -o  "src/main.o" "../src/main.cu"

что приводит к ошибке

ptxas fatal   : Unresolved extern function '_ZN11VirtualBaseD2Ev'
make: *** [src/main.o] Error 255

Я на Mac с установленной CUDA 7.5, но я попробовал то же самое на машине с Ubuntu 14.10 и CUDA 7.0, и получил те же результаты.


c++ inheritance cuda pure-virtual
person Dave    schedule 18.10.2015    source источник

Ответы (1)


arrow_upward
0
arrow_downward

После нескольких часов отладки, написания этого вопроса и наблюдения за ошибкой ptxas у меня возникло странное ощущение, что это деструктор базового класса, который не найден из-за D в конце _ZN11VirtualBaseD2Ev.

Я искал способы разобрать идентификатор, и действительно, D означает деструктор (а стандартный конструктор имеет C в том же месте).

Через несколько операторов отладки я понял, что когда Implementation impl; выходит за пределы области видимости, вызываются оба деструктора, сначала собственный, а затем базовые классы. Поскольку деструктор базового класса не имеет реализации, его нельзя вызвать, и возникает ошибка.

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

person Dave    schedule 18.10.2015
comment
И чтобы ответить на ваш первоначальный вопрос, да, вы можете использовать класс, который реализует виртуальные функции внутри ядра CUDA. Это можно сделать, создав экземпляры объектов этих классов на устройстве, а не на хосте. Ограничение состоит в том, что вы не можете передавать объекты таких классов с хоста на устройство. В ограничении не указано, что вы не можете использовать их на устройстве. - person Robert Crovella; 18.10.2015
comment
Базовый класс всегда должен предоставлять реализацию своего виртуального деструктора (например, virtual ~Base() = default;). Вы можете прочитать книгу по C++. - person user703016; 19.10.2015