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

【问题标题】：Can I override a CUDA host-and-device function with a host-only function?我可以使用仅主机功能覆盖 CUDA 主机和设备功能吗？
【发布时间】：2019-12-16 21:04:18
【问题描述】：

考虑以下程序：

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 告诉我，我不应该在覆盖方法时删除执行空间。我不是在问我自己的代码（我在这里没有引用），而是在问原理：

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

【问题讨论】：

标签： compiler-errors cuda overriding virtual-functions nvcc

【解决方案1】：

覆盖（虚拟）方法必须尊重被覆盖方法的执行空间选择。

“覆盖”仅与虚拟方法相关 - 因此您的 C::foo() 必须是标记为 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__-side 调用，而子类方法将应用于__host__-side 调用。但这也有点尴尬 - 我们需要在通过基类 ptr 作用于对象时调用 something。

【讨论】：