我可以使用仅主机功能覆盖 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 告诉我我不应该在覆盖方法时放弃执行 space。我问的不是我自己的代码（这里没有引用），而是原理：

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

Answer 1

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

"overriding" 仅与虚拟方法相关 - 因此您的 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

这个限制有意义吗？可以想象一种解释，其中 base-class 方法将应用于 __device__ 端调用，而 subclass 方法将应用于 __host__ 端调用。但这也有点尴尬——我们需要调用 something 当通过基 class ptr.

作用于一个对象时

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

Can I override a CUDA host-and-device function with a host-only function?

overriding

cuda

virtual-functions

compiler-errors

nvcc

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