gpu - Passing Host Function as a function pointer in __global__ OR __device__ function in CUDA

Question

I am currently developing a GPU version of a CPU function (e.g. function Calc(int a, int b, double* c, souble* d, CalcInvFunction GetInv )), in which a host function is passes as a function pointer(e.g. in above example GetInv is the host function of CalcInvFunction type). My question is, if i have to put Calc() function entirely in GPU, i have to pass the GetInv function as a function pointer argument in device function/kernel function, and is that possible?

Answer 1

Yes, for a GPU implementation of Calc, you should pass the GetInv as a __device__ function pointer.

It is possible, here are some worked examples:

Ex. 1

Ex. 2

Ex. 3

Most of the above examples demonstrate bringing the device function pointer all the way back to the host code. This may not be necessary for your particular case. But it should be fairly obvious from above how to grab a __device__ function pointer (in device code) and use it in a kernel.