这个问题很简单,但是让我先给你介绍一下我的框架。我有一个抽象类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__ 函数。注意:我删除的问题部分是错误的,请查看评论以了解原因。
__device__
函数可以是虚函数:http://docs.nvidia.com/cuda/cuda-c-programming-guide/#virtual-functions - m.s.__device__
函数可以是virtual
的。编程指南甚至提供了一些例子。我认为你所指的可能是,在主机上创建具有虚函数的对象无法传递到设备。 - Robert Crovella__host__
又是__device__
的虚函数。但你不能做的正是我所说的,即在主机上实例化带有虚函数的对象,然后将其传递到设备上。因为对象的虚函数表会被主机(函数)指针填充。当该类被复制到设备时,这些主机指针就变得毫无意义了。 - Robert Crovella__device__
函数可以是虚函数,这与你声称“我清楚地理解__device__
函数为什么不能是虚函数”的说法相矛盾。 - Michal Hosala