Handling loss of host functions in kernel in C

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi,

Still learning here. So I just discovered that host functions are inaccessible to the kernel. I’m going to split my question into two parts:

  1. First part lets stick to vanilla C only. For the program I am seeking to optimize the basic code (and its related data structures) have been given to me. A somewhat complex custom struct is used which requires a lot of helper functions for data access. Obviously once I move to the device I lose access to these, so for the moment I’ve had to resort to copy all the data I need for the calculation to a straight array and do the processing that way.

On the other hand I wouldn’t (think ?) it would be advisable to stuff a second copy of all your helper functions into the kernel (?!?).

Certainly there must be a better way to handle this scenario-- But I’m not sure what it is ?

  1. Let’s now pretend we are in C++ land where we have true objects that have methods. If I were to pass an entire object to the kernel, would any associated helper methods be pulled along with it (and thus usable) ? Or are we still in the same situation as above ?
2

Are you aware of the concept of __host__ __device__ functions?

1 Like
3

@njuffa No, I’ve only done the Fundamentals Course and at least off the top of my head don’t remember this coming up. Let me look into it. Thanks.

4

I might still ask though if you don’t mind-- What about scenario 2 ? Can you send a whole object to the kernel ?

5

yes, an object whose methods are decorated with __device__ will allow for those methods to be callable from device code. You can “create” an object in host code and use it in device code. You can also create objects in device code. There is a restriction in CUDA on passing objects with virtual methods between host and device code. There are numerous questions about using C++ objects with CUDA that you can find with a bit of searching.