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

我建议在设备代码中使用多态的方法是 在设备上创建对象。这将使用一组设备函数指针(而不是主机函数指针)设置vtable，并且诸如 this之类的问题证明了它可以工作。一阶近似，如果您有一种方法可以在宿主代码中创建一组多态对象，那么我不知道您为什么不能在设备代码中使用类似方法。问题实际上与互操作性有关-在主机和设备之间移动此类对象- 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。