具有虚拟继承传递到CUDA内核函数的对象

object with virtual inheritance passing to cuda kernel function

本文关键字：内核 CUDA 函数 对象 虚拟 继承      更新时间：2023-10-16

我能够将普通对象传递给内核函数作为副本。但是，当我在类层次结构中添加虚拟继承时，我收到了一个错误消息，说具有用户定义的复制构造函数的类不能用作内核启动的参数。但是，我没有用户定义的复印件。因此，我想这是因为虚拟继承的内部实现实现了某些不同类型的复制构建器。有人可以解释虚拟继承如何实际实现，并且是否有解决方法，或者在编写CUDA代码时绝对没有办法使用虚拟继承？

代码是这样的：

class Base {...};
class ChildA: public virtual Base {...};
class ChildB: public virtual Base {...};
class GrandChild: public ChildA, public ChildB {...};
__global__ void mykernel(Base x) {...}
int main() {
  GrandChild x;
  mykernel<<<1,1>>>(x);
  return 0;
}

编辑：这是我的猜测：我认为NVCC仅允许默认复制构建器，因为在这种情况下，它可以简单地使用cudamemcpyasync将参数推入设备内的呼叫堆栈。因此，它对编译器进行了硬编码，以便仅允许使用默认的复制构建器，但是具有虚拟继承的对象在内部具有不同类型的复制构建器，这触发了NVCC中的错误。但是，我仍然认为NVCC可以使用一种简单的方法，但前提是NVCC支持虚拟功能和其他高级C 功能。

感谢所有评论。根据链接@RobertCrovella提供的以下两个陈述，我不允许这样做。

  1. 不允许将参数作为参数传递给__global__函数一个具有虚拟函数的类的对象。
  2. 不允许作为参数传递给__global__函数一个从虚拟基类派生的类的对象。

与提供的链接@Joky一起，我相信原因是因为NVCC使用简单的内存副本将主机内存的参数从设备内存中的呼叫堆栈中推入呼叫堆栈。这就是为什么不允许非默认复制构建器的原因，因为只有默认的复制构造人与此简单的内存复制行为一致。

我正在使用的解决方案是打破一个继承关系并设置类型转换操作员，以保留缺少的链接以进行升级，如以下内容。这对我有用，因为所有这些类都是指针的包装器，并且类型转换的效率足够有效，即使价格是有时我必须进行明确的类型转换，例如以下调用mykernel2功能的情况。 

class Base {...};
class ChildA: public Base {...};
class ChildB: public Base {...};
class GrandChild: public ChildA {
public:
  operator ChildB () {...}
};
__global__ void mykernel(Base x) {...}
__global__ void mykernel2(ChildB y) {...}
int main() {
  GrandChild x;
  mykernel<<<1,1>>>(x);
  mykernel2<<<1,1>>>(ChildB(x));
  return 0;
}

相关文章：