c ++ - CUDA avec l'appel de méthodes virtuelles dans le noyau sur les objets qui a été créé sur l'hôte - Stack Overflow

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

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).

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

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).

Source