Link errors for struct/class member functions CUDA 2.1, MSVC 2008, XP x86

[Ojiisan]

We all know that a function can be compiled both for the device and the host simultaneously by CUDA to use in your host code. CUDA properly exports these symbols and you can easily use them in your program, be it C, or C++.

It seems, at least for me, that NVCC has a compiler bug where the same does NOT apply for struct/class member functions.

See the following snippet:

include.h

struct TEST {

  __device__ __host__ int callme();

};

int callme_global();

main_cu.cu

#include <cuda_runtime.h>

#include "include.h"

#pragma comment(lib, "cudart.lib")

int TEST::callme() {return 10;}

__device__ void dummy_kernel() {TEST me; me.callme();}

//__host__ void dummy_host() {TEST me; me.callme();} // XXX IMPORTANT!

int callme_global() {return 10;}();

main_cpp.cpp

#include <cuda_runtime.h>

#include "include.h"

int main() {

  TEST me;

  me.callme();

  callme_global();

}

Compiling with CUDA 2.1 and Visual Studio 2008 gives the following link error:

[b]1>main_cpp.obj : error LNK2019: unresolved external symbol “public: int __thiscall TEST::callme(void)” (?callme@TEST@@QAEHXZ) referenced in function _main

1>C:\cuda_test\Debug\cuda_test.exe : fatal error LNK1120: 1 unresolved externals[/b]

We see that TEST::callme() isn’t exported no matter what I do. The only “solution” I have found is writing a dummy host function in the .cu file (which is at the moment commented out), using those functions. Then, it seems NVCC sees that the function is used in host mode and generates the export for it; otherwise, even if host is specified in the declaration, the function is ignored.

It would be great if this issue could be fixed (together with class inheritance in emulation mode :) ). Or is this the same issue with no support for C++ in kernels?