CUDA内核函数可以作为虚函数吗？

Question

这个问题很简单，但是让我先给你介绍一下我的框架。我有一个抽象类AbstractScheme代表一种计算类型（一种方程的离散化，但这并不重要）。 每个实现都必须提供一个返回方案名称的方法，并实现一个受保护的函数，该函数是CUDA内核。基础抽象类提供了一个公共方法，该方法调用CUDA内核并返回内核完成所需的时间。

class AbstractScheme
{
public:
    /**
     * @return The name of the scheme is returned
     */
    virtual std::string name() const =0;

    /**
     * Copies the input to the device,
     * computes the number of blocks and threads,
     * launches the kernel,
     * copies the output to the host,
     * and measures the time to do all of this.
     *
     * @return The number of milliseconds to perform the whole operation
     *         is returned
     */
    double doComputation(const float* input, float* output, int nElements)
    {
        // Does a lot of things and calls this->kernel().
    }

protected:
    /**
     * CUDA kernel which does the computation.
     * Must be implemented.
     */
    virtual __global__ void kernel(const float*, float*, int) =0;
};

我还有这个基类的几个实现。但是当我尝试使用nvcc 7.0编译时，我会得到一个错误消息，这个错误消息涉及到我在AbstractScheme中定义函数kernel的那一行（如上列表中的最后一行）：

myfile.cu(60): error: illegal combination of memory qualifiers

我找不到任何资源表明内核不能是虚函数，但我有这种感觉是问题所在。你能解释一下背后的原理吗？我清楚地理解了为什么和如何 __device__ 函数不能是虚函数（虚函数是指向实际 [主机] 函数的指针存储在对象中，您无法从设备代码中调用这样的函数），但我不确定 __global__ 函数。注意：我删除的问题部分是错误的，请查看评论以了解原因。

- Spiros

---

__device__函数可以是虚函数：http://docs.nvidia.com/cuda/cuda-c-programming-guide/#virtual-functions - m.s.

__device__函数可以是virtual的。编程指南甚至提供了一些例子。我认为你所指的可能是，在主机上创建具有虚函数的对象无法传递到设备。 - Robert Crovella

你可以拥有一个既是 __host__ 又是 __device__ 的虚函数。但你不能做的正是我所说的，即在主机上实例化带有虚函数的对象，然后将其传递到设备上。因为对象的虚函数表会被主机（函数）指针填充。当该类被复制到设备时，这些主机指针就变得毫无意义了。 - Robert Crovella

你可以尝试使用普通的虚拟包装函数来调用内核，从而实现类似的功能。 - Robert Crovella

你不想编辑一下你的问题吗？已经有两个人告诉你__device__函数可以是虚函数，这与你声称“我清楚地理解__device__函数为什么不能是虚函数”的说法相矛盾。 - Michal Hosala

Answer 1

在CUDA对象模型中，无论是虚拟的还是非虚拟的，内核都不能成为类的成员。这就是编译错误的原因，即使编译器发出的错误消息并不特别明显。

- talonmies