I’ve been looking for example CUDA code that uses virtual functions, and so far I haven’t been able to find anything besides the snippet in section D.1.2 of the Programming Guide.
I’ve been working a bit with the example in section D.1.1, and have successfully compiled it and used the PixelRGBA class. However, if I modify it to incorporate inheritance and virtual functions, I get compiler errors.
For example:
class PixelRGBABase
{
public:
__device__ virtual unsigned char get_r()
{
return 3;
}
};
class PixelRGBA : public PixelRGBABase
{
public:
__device__ PixelRGBA(): r_(0), g_(0), b_(0), a_(0)
{ ; }
__device__ PixelRGBA(unsigned char r, unsigned char g, unsigned char b,
unsigned char a = 255): r_(r), g_(g), b_(b), a_(a)
{ ; }
__device__ unsigned char get_r()
{
return r_;
}
private:
unsigned char r_, g_, b_, a_;
friend PixelRGBA operator+(const PixelRGBA &,
const PixelRGBA &);
};
__device__
PixelRGBA operator+(const PixelRGBA & p1, const PixelRGBA & p2)
{
return PixelRGBA(p1.r_ + p2.r_,
p1.g_ + p2.g_,
p1.b_ + p2.b_,
p1.a_ + p2.a_);
The above code gives me the following error:
nvcc -ccbin /usr/bin -I. -I/usr/local/cuda/include --ptxas-options=-v -c pixelrgba.cu -o pixelrgba.cu.o
./pixelrgba.cu(20): Warning: Cannot tell what pointer points to, assuming global memory space
ptxas /tmp/tmpxft_00004313_00000000-2_pixelrgba.ptx, line 74; error : Feature 'functions as initial values' requires PTX ISA .version 2.1 or later
ptxas /tmp/tmpxft_00004313_00000000-2_pixelrgba.ptx, line 74; error : Feature 'functions as initial values' requires .target sm_20 or higher
ptxas /tmp/tmpxft_00004313_00000000-2_pixelrgba.ptx, line 75; error : Feature 'functions as initial values' requires PTX ISA .version 2.1 or later
ptxas /tmp/tmpxft_00004313_00000000-2_pixelrgba.ptx, line 75; error : Feature 'functions as initial values' requires .target sm_20 or higher
ptxas fatal : Ptx assembly aborted due to errors
make: *** [pixelrgba.cu.o] Error 255
I’m using a Tesla C2050 with driver version 270.41.06. According to the Programming Guide, virtual functions should work with GPU’s of compute capability 2.x.