See the following snippet (CUDA 2.1, MSVC2008, WindowsXP 32bits 8600GT):
class A {
public:
int a;
};
class B: public A {
public:
int b;
};
__global__ void aa() {
B bb;
bb.b = 6;
bb.a = 7;
}
[i][b]1>test.cu(13) : error C2039: ‘__b_1A’ : is not a member of ‘B’
1> test.cu(5) : see declaration of ‘B’
1>test.cu(13) : error C2228: left of ‘.a’ must have class/struct/union
[/b][/i]
When code is compiled for device mode, everything goes correctly. I had a look at nvcc internals (with -v) and I think I kinda figured out what is going wrong. Any device code is compiled to a source.gpu file, in C mode. In there you will see that class B has a member A __b__1A, and the assignment code is actually ((bb.__b__1A).a) = 7;. Now in emulation mode your code gets compiled in C++ mode so you keep the structures, but the global function is replaced by that in the .gpu file; which is C, and has the inherited class/struct renamed. Bamm, your host compiler stops working, and errors.
Any chance of this getting fixed any time soon? As it is now, I can only blind-test my device code which is really annoying :(
P.S. I presume in device mode the generated .c file is used in which the class is already transformed to the liking of the .gpu file; thus everything works.
P.P.S. Actually, another, more serious problem is with templates; again in emulation mode, device mode works. See the following:
template<class T> struct auto_any {
auto_any(T const &t): item(t) {}
T item;
};
template<class T> __device__ inline auto_any<T> contain(T const &t) {return t;}
__global__ void aa() {
int bb;
contain(bb);
}
[i][b]test.cu(14) : error C2526: ‘Z7containIiE8auto_anyIT_ERKS1’ : C linkage function cannot return C++ class ‘auto_any’ with [T=int]
test.cu(14) : error C2526: ‘Z7containIiE8auto_anyIT_ERKS1’ : C linkage function cannot return C++ class ‘auto_any’ with [T=int]
test.cu(14) : error C2512: ‘auto_any’ : no appropriate default constructor available with [T=int]
test.cu(14) : error C2562: ‘Z7containIiE8auto_anyIT_ERKS1’ : ‘void’ function returning a value
test.cu(14) : see declaration of '_Z7containIiE8auto_anyIT_ERKS1_'
[/b][/i]