我可以用仅主机功能覆盖CUDA主机和设备功能吗？

Question

考虑以下程序：

class A {
    __host__  __device__ void foo();
};

class B : A {
    __host__ void foo();
};

int main()
{
    A a; (void) a; 
    B b; (void) b;
}

此compiles（GodBolt）与nvcc 10。

但是，在更复杂的程序中，有时会出现以下错误（为了可读性而换行）：

whatever.hpp(88): error: execution space mismatch: overridden entity (function
"C::foo") is a __host__ __device__ function, but overriding entity (function "D::foo")
is a __host__ function

所以，nvcc告诉我，我[[not应该在覆盖方法时删除执行空间。我不是在问自己的代码（在这里没有引用），而是在问原理：

如果仅用__host__ __device__个函数（我认为是合理的）覆盖__host__个函数是可以接受的，那么nvcc怎么会出现这种错误？
• 或者，如果不允许，为什么上面的小程序不能编译？

Answer 1

重载虚拟方法必须尊重被重载方法的执行空间选择。

“ overriding”仅与虚拟方法有关-因此，在您的C::foo()必须为

marked virtual的情况下。实际上，如果我们在示例程序中将foo()标记为虚拟：

class A {
    virtual __host__  __device__ void foo();
};

class B : A {
    __host__ void foo(); // can say "override" here; but it doesn't matter
};

int main()
{
    A a; (void) a; 
    B b; (void) b;
}

这将是fail to compile：
<source>(6): error: member function declared with "override" does not override
a base class member

此限制有意义吗？可以想象一种解释，其中基类方法将应用于__device__端调用，而子类方法将应用于__host__端调用。但这也有点尴尬-当通过基类ptr对对象进行操作时，我们需要调用
something。