title: How to pass a function as a cuda kernel parameter?
body: <p>I would like to create generic cuda kernel that take a callable object as parameter (like lambda or function) and invoke it.</p>
  
  <p>I am having trouble to pass a device function to a cuda kernel as a parameter. </p>
  
  <p>I have cuda 9.2 with compute capability 3.5. I use gcc 9.3 on Debian 10. </p>
  
  <p>I tried this, compiled with <code>nvcc -arch=sm_35 --expt-extended-lambda main.cu -o test</code>:</p>
  
  <pre class="lang-cpp prettyprint-override"><code>    __host__ __device__ void say_hello()
      {
          printf("Hello World from function!\n");
      }
  
      template&lt;class Function&gt;
      __global__ void generic_kernel(Function f)
      {
          f();
      }
  
      int main() 
      {
              // this is working
          generic_kernel&lt;&lt;&lt;1,1&gt;&gt;&gt;([]__device__(){printf("Hello World from lambda!\n");});
          cudaDeviceSynchronize();
  
              // this is not working!
          generic_kernel&lt;&lt;&lt;1,1&gt;&gt;&gt;(say_hello); 
          cudaDeviceSynchronize();
  
          return 0;
      }
  </code></pre>
  
  <p>I expected to see both <code>Hello World from function!</code> and <code>Hello World from lambda!</code> but I only see the message from the lambda.</p>