CUDA 5.5 nvlink undefined reference (inheritance)

Question

I've been working on my GPU-raytracer implementation, but as I am new to CUDA I have some problems with compiling and linking the separate .cu files. My 2 classes: Shader and Lambert. Lambert inherits the interface Shader. When I compile I recieve the following errors:

Error 4 error MSB3721: The command ""G:\Development\CUDA Toolkit\CUDA Toolkit v5.5\bin\nvcc.exe"
-dlink -o "Debug\CUDA RayTracer.device-link.obj" -Xcompiler "/EHsc /W3 /nologo /Od /Zi /RTC1  
/MDd  " -L"P:\My Projects\CUDA Ray-Tracer\CUDA RayTracer\ThirdParty\SDL\lib\x86" -L"P:\My 
Projects\CUDA Ray-Tracer\CUDA RayTracer\CUDA RayTracer\\..\ThirdParty" -L"G:\Development\CUDA
Toolkit\CUDA Toolkit v5.5\lib\Win32" cudart.lib kernel32.lib user32.lib gdi32.lib winspool.lib
comdlg32.lib advapi32.lib shell32.lib ole32.lib oleaut32.lib uuid.lib odbc32.lib odbccp32.lib 
SDL.lib SDLmain.lib  -gencode=arch=compute_30,code=sm_30 -G --machine 32 Debug\Camera.cu.obj 
Debug\IShader.cu.obj Debug\Lambert.cu.obj Debug\Matrix.cu.obj Debug\Plane.cu.obj 
Debug\sdl.cu.obj Debug\cuda_renderer.cu.obj" exited with code -1.   
C:\Program Files(x86)\MSBuild\Microsoft.Cpp\v4.0\V110\BuildCustomizations\CUDA 5.5.targets  668


1>nvlink : error : Undefined reference to '_ZN6ShaderD1Ev' in 'Debug/IShader.cu.obj'
1>nvlink : error : Undefined reference to '_ZN6ShaderD0Ev' in 'Debug/IShader.cu.obj'
1>nvlink : error : Undefined reference to '_ZN6ShaderD2Ev' in 'Debug/Lambert.cu.obj'

I have no idea what '_ZN6ShaderD1Ev' means, I think everything is correct in my implementation (C++ wise, not sure what CUDA thinks about it). As far as I know CUDA 5.5 supports virtual functions and inheritance.

I have installed CUDA 5.5 Toolkit and I have enabled the "Generate Relocatable Device Code" in my Visual Studio 2012. Also I've set 'compute_30,sm_30'in order to use 'operator new' (my graphics card is capable of it - GTX670MX). My project consists only of .cu and .cuh files.

My source code:

//IShader.cuh
#ifndef I_SHADER_H
#define I_SHADER_H

#include "Vector3D.cuh"
#include "Color.cuh"
#include "IGeometry.cuh"

__device__ extern Vector cameraPos;
__device__ extern Vector lightPos;
__device__ extern Color lightColor;
__device__ extern float lightPower;
__device__ extern Color ambientLight;

class Shader
{
protected:
     Color _color;

public:
    __device__ Shader(const Color& color);

    __device__ virtual ~Shader();
    __device__ virtual Color shade(Ray ray, const IntersectionData& data) = 0;
};

#endif

//IShader.cu
#include "IShader.cuh"

__device__ Shader::Shader(const Color& color)
{
   this->_color = color;
}

// Lambert.cuh
#ifndef LAMBERT_H
#define LAMBERT_H

#include "IShader.cuh"

class Lambert : public Shader
{
public:
    __device__ Lambert(const Color& diffuseColor);
    __device__ Color shade(Ray ray, const IntersectionData& data);
};

#endif

//Lambert.cu
#include "Lambert.cuh"

Vector cameraPos;
Vector lightPos;
Color lightColor;
float lightPower;
Color ambientLight;

__device__ Lambert::Lambert(const Color& diffuseColor)
    : Shader(diffuseColor)
{
}

__device__ Color Lambert::shade(Ray ray, const IntersectionData& data)
{
    Color result = _color;

    result = result * lightColor * lightPower / (data.p - lightPos).lengthSqr();
    Vector lightDir = lightPos - data.p;
    lightDir.normalize();

    double cosTheta = dot(lightDir, data.normal);
    result = result * cosTheta;

    return result;
}

If you need more code I could give you link to github repo. I hope you can help me. Thanks in advance!

Answer 1

C++ enables different entities (e.g., functions) named with the same identifier to belong to different namespaces. To uniquely resolve names, the compiler uses name mangling, that is, it encodes additional information in the name of the involved entities. This is the reason why nvlink is referring to this "obscure" entity _ZN6ShaderD1Ev. In order to restore a more understandable name, a demangling operation is necessary.

Although demangling software exist, I'm often using an online demangler

c++filtjs

Using this page, you can discover that

_ZN6ShaderD1Ev

actually means

Shader::~Shader()

which, in turn, suggests that you are not defining the destructor for the Shader() class.