有什么方法可以将vtable从主机复制到设备(CUDA和C ++) [英] Is There Any Way To Copy vtable From Host To Device (CUDA & C++)

问题描述

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

It seems that Cuda does not allow me to "pass an object of a class derived from virtual base classes to __global__ function", for some reason related to "virtual table" or "virtual pointer".

我想知道有什么办法可以设置虚拟指针"吗?手动操作，以便我可以使用多态性?

I wonder is there some way for me to setup the "virtual pointer" manually, so that I can use the polymorphism?

推荐答案

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

I would like to provide a different way to fix the vtable which does not rely on copying the vtable between objects. The idea is to use placement new on the device to let the compiler generate the appropriate vtable. However, this approach also violates the restrictions stated in the programming guide.

#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;
}

这篇关于有什么方法可以将vtable从主机复制到设备(CUDA和C ++)的文章就介绍到这了，希望我们推荐的答案对大家有所帮助，也希望大家多多支持IT屋！