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

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

本文关键字：函数 实现 CUDA 内核 我可以      更新时间：2023-10-16

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

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

我知道

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

，因为虚表在主机上创建并复制到GPU时将是垃圾。但是我没有将对象传递给内核，而是在内核内部构造对象，这应该不会导致虚函数表问题。

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的机器上尝试了同样的事情，产生了相同的结果。

经过几个小时的调试，编写这个问题并盯着ptxas错误，我有一种奇怪的感觉，这是没有找到的基类的析构函数，因为D接近_ZN11VirtualBaseD2Ev的末尾。

我寻找了一些方法来修改标识符，实际上，D代表析构函数(标准构造函数在同一位置有一个C)。

几个调试语句之后，我意识到，当Implementation impl;超出作用域时，两个析构函数都被调用，它自己的第一个，然后是基类。由于基类的析构函数没有实现，因此不能调用它，并抛出错误。

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

相关文章：