CODEWITHSUNDEEP
c++cuda
Arjun K R
Arjun K R

Reputation: 462

Passing a class object to a kernel

Can we pass an object to a kernel function ?

Consider I have a class 

class MyClass
{
    public:
        int value;
        float rate;
        MyClass()
        {
            value = 0; rate = 0;
        }
        MyClass(int v,float r)
        {
            value = v; rate = r;
        }
};

and my Kernel takes an array of the objects of MyClass

__global__ void MyKernel(MyClass * mc)
{
    //Some Calculation
}

Can I pass the array ? How to allocate memory ?? Right now I tried with the following code got CudaMemcpy error

cudaError_t cudaStatus;

MyClass darr[10] ;
cudaStatus = cudaMalloc((void**)&darr, size * sizeof(MyClass));

if (cudaStatus != cudaSuccess) {
    fprintf(stderr, "cudaMalloc failed!");
    goto label1;
}

cudaStatus = cudaMemcpy(darr, arr, size * sizeof(MyClass), cudaMemcpyHostToDevice);
//arr is a host array

Upvotes: 10

Views: 11229

Answers (1)

talonmies
talonmies

Reputation: 72349

There are a few problems here, not all directly related to whatever error you are seeing.

Firstly, you will have to define each class method in both the host and device so that the class can be instantiated in both memory spaces (when you do a copy, only the data members of each instance are copied). So your class declaration should look something like this:

class MyClass
{
    public :
        int value;
        float rate;
        __device__ __host__ MyClass()
        {
            value = 0; rate = 0;
        }
        __device__ __host__ MyClass(int v,float r)
        {
            value = v; rate = r;
        }
        __device__ __host__ ~MyClass() {};
}

You then need to correctly allocate the device memory. If you want an array of MyClass with 10 members on the device, allocate and copy it to the device like this:

MyClass arr[10];
MyClass *darr;
const size_t sz = size_t(10) * sizeof(MyClass);
cudaMalloc((void**)&darr, sz);
cudaMemcpy(darr, &arr[0], sz, cudaMemcpyHostToDevice);

[disclaimer: all code written in browser, never complied or tested, use at own risk]

You can then pass darr as an argument to a kernel from the host.

Upvotes: 18

Related Questions