我正在努力解决一个似乎有些晦涩的问题。
我正在一个框架上工作,在该框架上,用户可以提供抽象基类的实现,在进行了一些魔术和代码生成步骤之后,该抽象基类将在CUDA内核中使用。
我知道
“不允许将带有虚函数的类的对象作为参数传递给全局函数。”
因为在主机上创建vtable并将其复制到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的计算机上尝试了相同的操作,产生了相同的结果。
更多的时间调试,写了这个问题,并在ptxas错误盯着后,我有奇怪的感觉,这是没有找到基类的析构函数,因为D
接近年底_ZN11VirtualBaseD2Ev
。
我寻找了拆解标识符的方法,实际上,还D
寻找了析构函数的立场(标准构造函数C
在同一位置有一个)。
后来的一些调试语句中,我意识到,当Implementation impl;
超出范围时,两个析构函数都被称为,它是第一个,然后是基类。由于基类的析构函数没有实现,因此无法调用它,并引发错误。
编辑:此析构函数调用当然不是CUDA问题,而是标准C ++例程。此外,正如罗伯特·克罗维拉(Robert Crovella)在评论中指出的那样,CUDA确实支持在设备上实例化实现虚拟功能的类。
本文收集自互联网,转载请注明来源。
如有侵权,请联系[email protected] 删除。
我来说两句