The following code fails with “nvcc --device-emulation -c test.cu”
test.cu:14: error: ‘_ZZN8MyStruct5helloEvEs’ was not declared in this scope
#include <stdio.h>
#include <cuda.h>
struct MyStruct {
__device__ void hello();
};
__device__ void MyStruct::hello() {
printf("hello world.\n");
}
__global__ void device_main() {
MyStruct s;
s.hello();
}
int cuda_main() {
device_main<<< 1, 1 >>>();
return 0;
}
changing the second block to the following fixes the problem
__device__ void globalscope_hello() {
printf("hello world.\n");
}
__device__ void MyStruct::hello() {
globalscope_hello();
}
nvcc “release 2.0, V0.2.1221”, “built on Thu_Jun_19_04:48:21_PDT_2008”
cheers,
Nicholas
Actually, it seems worse than that; it seems to fail with any pointer if it’s called from a device class function.
__device__ void globalscope_hello(const char *str) {
printf("hello world '%s'.\n", str);
}
__device__ void MyStruct::hello() {
globalscope_hello("something");
}
getting rid of the class works
__device__ __host__ void globalscope_hello(const char *str) {
printf("hello world '%s'.\n", str);
}
__global__ void device_main() {
globalscope_hello("asdfq");
}
any workarounds would be much appreciated. I’d prefer to be able to use classes and debug… unfortunate this doesn’t seem to work immediately with cuda 2.0.
thanks,
Nicholas
C++ is unsupported–some features may work sporadically, as you’ve seen, but things like classes and operator overloading will almost certainly break from version to version.