CODEWITHSUNDEEP
c++classcudavirtual
waskithebest
waskithebest

Reputation: 95

CUDA with invoking virtual methods in kernel on objects that was created on Host

Description of my problem is in this code below.

class A
{
public :
        __device__ __host__ virtual void foo() = 0;
        __device__ __host__ void foo2()
        {
                //something to do...
        }
};

class B : public A
{
public :
        __device__ __host__ B()
        { }
        //implementing virtual method foo from A
        void foo()
        {
                //something to do...
        }
};

__global__ void Test(A* ptr_device)
{
        ptr_device->foo2(); //it's okey
        ptr_device->foo(); //ERROR !!!
}

void main()
{
        B* b_host = new B();

        A* a_host = (A*)malloc(sizeof(A));
        memcpy(a_host, b_host, sizeof(A));

        a_host->foo();
        a_host->foo2();
        //it's okey

        A* a_device;
        cudaMalloc((void**)&a_device, sizeof(A));
        cudaMemcpy(a_device, a_host, sizeof(A), cudaMemcpyHostToDevice);
        Test<<<1, 1>>>(a_device);
}

The problem is when I want to invoke a virtual method of class A on kernel. First I'm creating a object B, then mallocing space for A object and copying B to A. On Host code it works fine but in Device code it gets an error in invoking foo method

Upvotes: 0

Views: 758

Answers (1)

PlasmaHH
PlasmaHH

Reputation: 16046

You can not just memcpy non-POD objects around and expect them to work, you invoke quite some UB here.

What happens in your case is likely that the object has some virtual table pointer that gets copied over, but has no meaning within that other execution environment as it almost certainly will not point to the virtual table that exists there (given it there is even one).

Upvotes: 1

Related Questions