有什么方法可以将vtable从主机复制到设备CUDA和C ++

问题描述

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

我想知道是否可以通过某种方式手动设置“虚拟指针”，以便可以使用多态性？

解决方法

有什么方法可以将vtable从主机复制到设备

您不想将vtable从主机复制到设备。主机上的vtable（即在主机上创建的对象中）在vtable中具有一组主机函数指针。将此类对象复制到设备时，vtable不会被更改或“固定”，因此最终在设备上出现一个对象，该对象的vtable充满了主机指针。

如果您随后尝试调用这些虚拟功能之一（使用设备代码中的设备上的对象），则会发生不良情况。 vtable中列出的数字函数入口点是在设备代码中没有任何意义的地址。

以便我可以使用多态性

我建议在设备代码中使用多态的方法是在设备上创建对象。这使用一组设备函数指针（而不是主机函数指针）设置了vtable，并且诸如this之类的问题证明了它可以工作。一阶近似，如果您有一种方法可以在宿主代码中创建一组多态对象，那么我不知道您为什么不能在设备代码中使用类似方法。问题确实与互操作性有关-在主机和设备之间移动此类对象-the stated limitations in the programming guide所指的是这种情况。

我想知道是否可以通过某种方式手动设置“虚拟指针”

可能有。为了共享知识，我将概述一种方法。但是，我对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。

,

我想提供一种不同的方法来修复vtable，它不依赖于在对象之间复制vtable。这个想法是在设备上使用new放置，以使编译器生成适当的vtable。但是，这种方法也违反了编程指南中规定的限制。

#include <cstdio>

struct A{
    __host__ __device__
    virtual void foo(){
        printf("A\n");
    }
};

struct B : public A{

    B(int i = 13) : data(i){}

    __host__ __device__
    virtual void foo() override{
        printf("B %d\n",data);
    }

    int data;
};

template<class T>
__global__
void fixKernel(T* ptr){
    T tmp(*ptr);

    new (ptr) T(tmp);
}

__global__
void useKernel(A* ptr){
    ptr->foo();
}


int main(){

    A a;
    a.foo();

    B b(7); 
    b.foo();

    A* ab = new B();

    ab->foo();

    A* d_a;
    cudaMalloc(&d_a,sizeof(A));
    cudaMemcpy(d_a,&a,sizeof(A),cudaMemcpyHostToDevice);

    B* d_b;
    cudaMalloc(&d_b,sizeof(B));
    cudaMemcpy(d_b,&b,sizeof(B),cudaMemcpyHostToDevice);

    fixKernel<<<1,1>>>(d_a);

    useKernel<<<1,1>>>(d_a);

    fixKernel<<<1,1>>>(d_b);

    useKernel<<<1,1>>>(d_b);

    cudaDeviceSynchronize();

    cudaFree(d_b);
    cudaFree(d_a);
    delete ab;
}