How to pass a pointer to a device function as an argument to a kernel function?

Question

This question is about CUDA C/C++ programming. I tried to search for it quite a lot but didn't find appropriate question and therefore neither answer too.
I have 1 device function, 1 kernel function and the main function:

typedef float (*pfunc)(float arg);

__device__ float dev_func(float arg) {
    return arg * arg;
}

__global__ void ker_func(pfunc fnc) {
    printf("%f\n", fnc(2));
}

int main(void) {
    pfunc fnc = dev_func;
    //now how do I copy this pointer to device memory?
    ker_func<<<1,1>>>(...);
    return 0;
}

Answer 1

From the CUDA programming guide:

The address of a __global__ function taken in host code cannot be used in device code (e.g. to launch the kernel). Similarly, the address of a __global__ function taken in device code cannot be used in host code.

It is not allowed to take the address of a __device__ function in host code.

So you've two options:

Define your __device__ function pointer globally and call it in kernel.

typedef float (*pfunc)(float arg);

__device__ float dev_func(float arg) {
    return arg * arg;
}

// create device function pointer here
__device__ pfunc dev_func_ptr = dev_func;

__global__ void ker_func() {
    // call function through device function pointer
    printf("%f\n", dev_func_ptr(2));
}

If you want to pass function pointer to kernel as argument, then:

#define gpuErrchk(val) \
    cudaErrorCheck(val, __FILE__, __LINE__, true)
void cudaErrorCheck(cudaError_t err, char* file, int line, bool abort)
{
    if(err != cudaSuccess)
    {
        printf("%s %s %d\n", cudaGetErrorString(err), file, line);
        if(abort) exit(-1);
    }
}

typedef float (*pfunc)(float arg);

__device__ float dev_func(float arg) {
    return arg * arg;
}

// create device function pointer here
__device__ pfunc dev_func_ptr = dev_func;

__global__ void ker_func(pfunc fnc) {
    // call function through device function pointer
    printf("%f\n", fnc(2));
}


int main(int argc, char** argv)
{
    // create a host function pointer
    pfunc host_function_ptr;
    // copy function pointer value from device to host
    gpuErrchk(cudaMemcpyFromSymbol(&host_function_ptr, dev_func_ptr, sizeof(pfunc)));
    // pass the copied function pointer in kernel
    ker_func<<<1,1>>>(host_function_ptr);

    gpuErrchk(cudaPeekAtLastError());
    gpuErrchk(cudaDeviceSynchronize());

    return 0;
}