如何在CUDA代码中使用虚函数(C++)

最新推荐文章于 2022-09-19 12:56:00 发布

Pxy

最新推荐文章于 2022-09-19 12:56:00 发布

阅读量1.4k

点赞数 3

分类专栏： cuda编程文章标签： cuda c语言可扩展

本文链接：https://blog.csdn.net/u010445006/article/details/78219033

版权

本文介绍了在CUDA代码中使用虚函数以实现可扩展性的方法。关键点在于对象需要在设备上创建。错误示例中，由于在设备上未正确初始化指向虚函数表的指针导致程序崩溃。修正方案包括在设备内存中动态分配对象，并使用二重指针以避免参数复制的问题。最后，强调了在设备上正确释放内存的重要性。

摘要由CSDN通过智能技术生成

最近做的工作需要在CUDA上使用虚函数来减轻工作量，增加软件的可扩展性。我原以为按照CPU上的编程方法就可以正常使用，然而发现自己果真是too young, too naive了。在查找了很多材料后才找到解决方案，因而记录下来分享给大家。

简介

想要在device上使用虚函数，最重要的一点：

The objects simply need to be created on the device.

上面这句话的意思就是类的实例需要在device上被创建。
换言之，不能使用普通的cudaMalloc方法，而应该在__device__或__global__函数里使用class *A=new A;语句。

举例

如果上面解释的不清楚，请看这个例子，该例子仅仅解释我的解决方案，在实际应用中，永远不要使用带有虚函数的类的数组，否则你会哭的。
1.定义含有虚函数的父类Base和子类Derived，定义方法与CPU上的无二致。

class Base
{
public:
    __host__ __device__ Base() {}
    __host__ __device__ ~Base() {}

    __host__ __device__ virtual bool fun(int index)
    {
        printf("No.%d Base.fun(int)\n", index);
        return false;
    }

    int value;
};

class Derived :public Base
{
public:
    __host__ __device__ Derived() {}
    __host__ __device__ ~Derived() {}
    __host__ __device__ virtual bool fun(int index)
    {
        printf("No.%d Derived.fun(int)\n", index);
        return true;
    }
};

接下来的步骤是我们平时使用的方法，但请注意：它是错的。在这里举一个错误的例子来对CUDA中使用虚函数进行更清楚地讲解。
2.定义一个名为TestFun的__global__函数，传入一个指向Base数组的指针，对数组中的每个元素调用虚函数fun，我们所期待的是如果传入的是Derived数组的指针，则调用Derived的fun。

__global__ void TestFun(bool *d_result, Base *d_derives, size_t size)
{
    int myId=blockDim.x*blockIdx.x+threadIdx.x;
    if (myId >= size)  
        return;