CUDA 内核可以是虚函数吗？答案

【问题标题】：Can CUDA kernels be virtual functions?CUDA 内核可以是虚函数吗？
【发布时间】：2015-06-29 16:18:58
【问题描述】：

这个问题非常直截了当，但让我概述一下我的框架。我有一个抽象类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__ 的功能。

编辑：我删除的部分问题是错误的。请查看 cmets 以了解原因。

【问题讨论】：

__device__ 函数可以是虚拟的：docs.nvidia.com/cuda/cuda-c-programming-guide/…
__device__ 函数可以是virtual。编程指南甚至gives some examples。我想您所指的可能是在主机上创建的具有virtual 函数cannot be passed to the device 的对象。
你可以同时拥有__host__和__device__的虚函数，你不能做的正是我说的，在宿主机上用虚函数实例化一个对象，然后传递给装置。因为对象虚函数表被主机（函数）指针填充。当类被复制到设备时，那些主机指针就没有意义了。
你可以用调用内核的普通虚拟包装函数来完成类似的事情。
您不想编辑您的问题吗？两个人已经告诉过你__device__ 函数可以是virtual。这与您声称“我清楚地理解__device__ 函数不能是虚函数的方式和原因”形成鲜明对比。

标签： c++ cuda polymorphism virtual-functions

【解决方案1】：

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

【讨论】：