title: Call device function from global function in pycuda
body: <p>I'm newbie to PyCUDA. I want to call function declared with <code>__device__</code> from function declared with <code>__global__</code>. How can I do this in pyCUDA?  </p>
  
  <pre><code>import pycuda.driver as cuda  
  from pycuda.compiler import SourceModule  
  import numpy as n  
  import pycuda.autoinit  
  import pycuda.gpuarray as gp
  
  d=gp.zeros(shape=(128,128),dtype=n.int32)  
  h=n.zeros(shape=(128,128),dtype=n.int32)  
  mod=SourceModule("""  
        __global__ void  matAdd(int *a)  
      {  
              int px=blockIdx.x*blockDim.x+threadIdx.x;  
              int py=blockIdx.y*blockDim.y+threadIdx.y;         
              a[px*128+py]+=1;   
              matMul(px);
  
      }  
        __device__ void matMul( int px)
      {
        px=5;
      }  
  
  """)
  
  m=mod.get_function("matAdd")  
  m(d,block=(32,32,1),grid=(4,4))  
  d.get(h)  
  </code></pre>
  
  <p>Above code is giving me following error</p>
  
  <pre><code>7-linux-i686.egg/pycuda/../include/pycuda kernel.cu]  
  [stderr:  
  kernel.cu(8): error: identifier "matMul" is undefined  
  
  kernel.cu(12): warning: parameter "px" was set but never used  
  
  1 error detected in the compilation of "/tmp/tmpxft_00002286_00000000-6_kernel.cpp1.ii".  
  ]  
  </code></pre>