Copy an object to device?

Question

Can I copy a C++ object to the device?

say I have:

class CudaClass
{
public:
int* data;
CudaClass(int x) {
    data = new int[1]; data[0] = x;
}
};

__global__ void useClass(CudaClass cudaClass)
{
    printf("%d" cudaClass.data[0]);
};


int main()
{
    CudaClass c(1);
}

Now how do I copy "c" to device memory and launch kernel "useClass"?

Answer 1

Yes, you can copy an object to the device for use on the device. When the object has embedded pointers to dynamically allocated regions, the process requires some extra steps.

See my answer here for a discussion of what is involved. That answer also has a few samples code answers linked to it.

Also, in your class definition, if you want certain functions to be usable on the device, you should decorate those functions appropriately (i.e. probably with __device__ __host__);

EDIT: In response to a question (now deleted) here is the simplest sample code I could come up with based on the supplied code:

#include <stdio.h>

class CudaClass
{
public:
int* data;
CudaClass(int x) {
    data = new int[1]; data[0] = x;
}
};

__global__ void useClass(CudaClass *cudaClass)
{
    printf("%d\n", cudaClass->data[0]);
};




int main()
{
    CudaClass c(1);
    // create class storage on device and copy top level class
    CudaClass *d_c;
    cudaMalloc((void **)&d_c, sizeof(CudaClass));
    cudaMemcpy(d_c, &c, sizeof(CudaClass), cudaMemcpyHostToDevice);
    // make an allocated region on device for use by pointer in class
    int *hostdata;
    cudaMalloc((void **)&hostdata, sizeof(int));
    cudaMemcpy(hostdata, c.data, sizeof(int), cudaMemcpyHostToDevice);
    // copy pointer to allocated device storage to device class
    cudaMemcpy(&(d_c->data), &hostdata, sizeof(int *), cudaMemcpyHostToDevice);
    useClass<<<1,1>>>(d_c);
    cudaDeviceSynchronize();
    return 0;
}

In the interest of brevity/clarity I have dispensed with the usual cuda error checking.

Responding to the question, you cannot allocate storage directly from the host using the pointer in the device-based class. This is because cudaMalloc expects an ordinary host based pointer storage, such as what you get with:

int *hostdata;

cudaMalloc cannot work with a pointer whose storage is already on the device. This will not work:

cudaMalloc(&(d_c->data), sizeof(int));

because it requires dereferencing a device pointer (d_c) in host code, which is not allowed.