title: cuda device function and templates
body: <p>I am using CUDA 7 and am trying to pass a function as a template parameter to a device function as follows:</p>
  
  <pre><code>typedef float(*Op)(float, float);
  
  template&lt;typename Op&gt;
  __device__ bool is_maxima(float ax, float ay, cudaTextureObject_t current)
  {
      // I try to use the passed function as:
      float cv = tex2D&lt;float&gt;(current, ax, ay);    
      float pv = tex2D&lt;float&gt;(current, ax - 1.f, ay);
      if (Op(cv, pv) != cv) return false;
      return true;
  }
  </code></pre>
  
  <p>Now from my <code>global</code> cuda kernel, I call this as follows:</p>
  
  <pre><code>__global__ void detect_keypoints(cudaTextureObject_t current
                                   float *result, int width, int height)
  {
      const int x = __mul24(blockDim.x, blockIdx.x) + threadIdx.x;
      const int y = __mul24(blockDim.y, blockIdx.y) + threadIdx.y;
  
      float c = tex2D&lt;float&gt;(current, ax + 0.5f, ay + 0.5f);
      float ax = x + 1; float av = y + 1;
  
      if (is_maxima&lt;fmaxf&gt;(ax + 0.5f, ay + 0.5f, current))
          result[y * width + height] = 1.f;
  }
  </code></pre>
  
  <p>However, it gives me a compiler error saying:</p>
  
  <pre><code>error: no instance of function template "is_maxima" matches the argument 
  list
          argument types are: (float, float, cudaTextureObject_t)
  </code></pre>
  
  <p>Is passing of functions as template parameters not allowed in CUDA device functions? I was under the impression that CUDA now supports all of C++ features.</p>
  
  <p>For completeness, <code>fmaxf</code> is defined in the CUDA SDK as:</p>
  
  <pre><code>inline float fmaxf(float a, float b)
  {
      return a &gt; b ? a : b;
  }
  </code></pre>