c++ - 有什么方法可以将表从主机复制到设备(CUDA 和 C++)

Question

由于某些与“虚拟表”或“虚拟指针”相关的原因，Cuda 似乎不允许我“将从虚拟基类派生的类的对象传递给 __global__ 函数” .

我想知道有没有办法让我手动设置“虚拟指针”，以便我可以使用多态性？

Answer 1

Is There Any Way To Copy vtable From Host To Device

您不会想将 vtable 从主机复制到设备。主机上的 vtable(即在主机上创建的对象中)在 vtable 中有一组主机函数指针。当您将此类对象复制到设备时，vtable 不会被更改或“修复”，因此您最终会在设备上得到一个对象，其 vtable 充满了主机指针。

如果您随后尝试调用这些虚函数之一(使用设备上的对象，从设备代码)，就会发生不好的事情。 vtable 中列出的数字函数入口点是在设备代码中没有任何意义的地址。

so that I can use the polymorphism

对于在设备代码中使用多态性的方法，我的建议是在设备上创建对象。这会使用一组设备函数指针而不是主机函数指针和诸如 this 之类的问题来设置 vtable。证明它有效。对于一阶近似，如果您有办法在主机代码中创建一组多态对象，我不知道为什么您不能在设备代码中使用类似的方法。这个问题确实与互操作性有关 - 在主机和设备之间移动此类对象 - 这就是 the stated limitations in the programming guide指的是。

I wonder is there some way for me to setup the "virtual pointer" manully

可能有。为了分享知识，我将概述一种方法。但是，我对 C++ 的了解还不足以说明这是否可以接受/合法。我唯一能说的是在我非常有限的测试中，它似乎有效。 但我认为它是不合法的，因此我不建议您将此方法用于实验以外的任何用途。即使我们不解决它是否合法，也已经声明CUDA 限制(如上所述)您不应尝试在主机和设备之间传递具有虚拟功能的对象。因此，我仅将其作为观察提供，这可能对实验或研究很有趣。我不建议将它用于生产代码。

this thread 中概述了基本思想.它基于这样的想法，即普通对象拷贝似乎不会复制虚函数指针表，这对我来说很有意义，但对象作为一个整体确实包含该表。因此，如果我们使用这样的方法:

template<typename T>
__device__ void fixVirtualPointers(T *other) {
        T temp =  T(*other); // object-copy moves the "guts" of the object w/o changing vtable
        memcpy(other, &temp, sizeof(T)); // pointer copy seems to move vtable
}

似乎可以获取给定对象，创建该类型的新“虚拟”对象，然后通过对象的基于指针的拷贝“修复”vtable(考虑整个对象的大小)而不是“典型的”对象拷贝。使用它需要您自担风险。 This blog阅读起来可能也很有趣，尽管我不能保证其中任何陈述的正确性。

除此之外，这里还有关于cuda标签的各种其他建议，不妨回顾一下them .