我可以使用在 CUDA 内核中实现纯虚函数的类吗？答案

【问题标题】：Can I use a class that implements pure virtual functions inside a CUDA kernel?我可以使用在 CUDA 内核中实现纯虚函数的类吗？
【发布时间】：2016-01-16 19:32:50
【问题描述】：

我正在努力解决一个看起来有点模糊的问题。

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

我知道

“不允许将具有虚函数的类的对象作为参数传递给全局函数。”

因为 vtable 在主机上创建然后复制到 GPU 时将是垃圾。但是我没有将对象传递给内核，我在内核内部构造了对象，这应该不会导致 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 的机器上尝试了相同的操作，得到了相同的结果。

【问题讨论】：

标签： c++ inheritance cuda pure-virtual

【解决方案1】：

经过更多小时的调试，写了这个问题，盯着ptxas错误，我有一种奇怪的感觉，这是没有找到的基类的析构函数，因为D接近尾端的@987654324 @。

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

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

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

【讨论】：

回答你原来的问题，是的，你可以使用一个在 CUDA 内核中实现虚函数的类。这样做的方法是在设备上实例化这些类的对象，而不是在主机上。限制是您不能将此类类的对象从主机传递到设备。该限制并未说明您无法在设备上使用它们。
基类应始终提供其虚拟析构函数的实现（例如virtual ~Base() = default;）。你可能想read a C++ book。