passing C++ objects to CUDA kernel

Question

I have a piece of code which consists of objects which I do not have the permission to modify. On the other hands I have to write CUDA kernel for doing some computations. At the moment I am converting the required information into an array or struct and passing it. Is it possible to pass objects as such to CUDA kernel and is there any way for CUDA kernel to access private members of the class. Some people have also suggested thrust to me.

Answer 1

You can use c++ functionality in your CUDA code. It has some additional requirements. For example kernel function can't me a member of class. For usage you own data structures it need to fit some requirements.

• Members that you need to use in device code should be annotated as __device__
• You need to implement __device__ constructors.
• If your structure encapsulate pointer this need to be allocated in device memory.
• You can use templates, but you can't use classes with virtual functions.

You can see some examples of c++ data structures in programming guide Appendix D. C/C++ Language Support.

As for your concrete questions:

1. Yes, you can pass objects in kernel. You need to implement __device__ copy-constructor for it.
2. You can access private members in class methods. You should implement __device__ accessor method to use it outside class. But you can't pass class with private member through kernel call.