It is not possible to naively use a __device__-function's pointer in host-side code. However, it is possible to use it if you copy its address from a global device-side variable which holds the address, like so:
__device__ void foo(int x) { /*... */ }
__device__ void (*ptr_on_device) (int x) = foo;
void bar() {
void (*ptr_on_host)(int x);
cuda::memory::region_t pod_region = cuda::memory::locate(ptr_on_device);
cuda::memory::copy(&ptr_on_host, pod_region.data(), pod_region.size());
// or using raw CUDA API calls:
// cudaMemcpyFromSymbol(&ptr_on_host, ptr_on_device, sizeof(void (*)(int)));
// ... and check the error
// Now do stuff with ptr_on_host
}
Perhaps we could add a mechanism abstracting the above to the library. While it would be host-side code mostly, passing device-side function pointers to kernels is definitely a useful "tool" for kernel authors to have.
It is not possible to naively use a
__device__
-function's pointer in host-side code. However, it is possible to use it if you copy its address from a global device-side variable which holds the address, like so:See also this question on StackOverflow.
Perhaps we could add a mechanism abstracting the above to the library. While it would be host-side code mostly, passing device-side function pointers to kernels is definitely a useful "tool" for kernel authors to have.