我可以使用 class 在 CUDA 内核中实现纯虚函数吗？

Question

我正在努力解决一个似乎有点晦涩的问题。

我正在开发一个框架，用户可以在该框架中提供抽象基础的实现 class，经过一些魔术和代码生成步骤后，它将在 CUDA 内核中使用。

我知道

"It is not allowed to pass as an argument to a global function an object of a class with virtual functions. "

因为在主机上创建然后复制到 GPU 时 vtable 将是垃圾。但是我没有把对象传递给内核，我在内核内部构造对象，应该不会引起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;
}

代码是用 Nsights 自动生成的 makefile 编译的

/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

我在安装了 CUDA 7.5 的 Mac 上，但我在装有 Ubuntu 14.10 和 CUDA 7.0 的机器上尝试了同样的操作，得到了相同的结果。

Answer 1

经过几个小时的调试，写出这个问题，盯着ptxas错误，我有一种奇怪的感觉，这是找不到的基class的析构函数，因为D _ZN11VirtualBaseD2Ev.

接近尾声

I looked for ways to demangle the identifier 实际上，D 代表析构函数（标准构造函数在同一位置有一个 C）。

几个调试语句之后，我意识到，当 Implementation impl; 超出范围时，两个析构函数都被调用 ，首先是自己的，然后是基 class之后。由于baseclass'的析构函数没有实现，所以无法调用，报错

编辑：这个析构函数调用当然不是 CUDA 问题，而是标准 C++ 例程。此外，正如 Robert Crovella 在评论中指出的那样，CUDA 确实支持 classes，如果它们在设备上实例化，则实现虚函数。

我可以使用 class 在 CUDA 内核中实现纯虚函数吗？

Can I use a class that implements pure virtual functions inside a CUDA kernel?

c++

inheritance

cuda

pure-virtual