Warp Invalid PC, device function pointer

Hey. I asked this question, but it looks like I am not going to get an answer there. So I am coming to you for help
Thank you, Roni

The comment I put on your question there is the direction you’ll need to follow. If your question is actually “what am I doing wrong?” then that comment gives the explanation for what you are doing wrong, in boldface.

In the link you sand I found
you still cannot pass nvstd::function objects initialized in host code to device code (and vice versa).. But it dose not say how to do it

Use a functor instead. My comment said that also.


$ cat t478.cu
#include "nvfunctional"
#include <stdio.h>
template<typename Func>
__global__ void testKernal(size_t size, Func func) {
    for(int i = 0;i < size; i++) {
        printf("func(%d) = %ld\n", i, func(i));
class myincfunc {

  __host__ __device__
  size_t operator()(const size_t val) {
   return val+1;

class MyClass {
    MyClass() {
        myincfunc func;
        testKernal<<<1,1>>>(10, func);
        cudaError_t err = cudaDeviceSynchronize();
        if(cudaSuccess != err) {
            printf("error with test kernal: %d\n", err);

int main(){

  MyClass a;
$ nvcc -std=c++11 -o t478 t478.cu
$ ./t478
func(0) = 1
func(1) = 2
func(2) = 3
func(3) = 4
func(4) = 5
func(5) = 6
func(6) = 7
func(7) = 8
func(8) = 9
func(9) = 10

roughly speaking, functor = “function wrapped in an object”

you cannot pass a bare function (or function pointer), not even a nvstd::function, from host to device, safely, unless you take care to capture the function pointer in device code. Ordinary definition or capture in host code will not work. There are many questions on stackoverflow that cover these concepts.

Here is a link to many more examples/resources:


Thank you for the grate response!
In my real case (not the example I gavie) I need 2 functions, one inc and one dec for example. I can use 2 functors but is there another way? I read something about cudamemcpy/cudamemcpysymble, im going to check it out when I get home.