I am using CUDA 7 and am trying to pass a function as a template parameter to a device function as follows:

```
typedef float(*Op)(float, float);
template<typename Op>
__device__ bool is_maxima(float ax, float ay, cudaTextureObject_t current)
{
// I try to use the passed function as:
float cv = tex2D<float>(current, ax, ay);
float pv = tex2D<float>(current, ax - 1.f, ay);
if (Op(cv, pv) != cv) return false;
return true;
}
```

Now from my `global`

cuda kernel, I call this as follows:

```
__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<float>(current, ax + 0.5f, ay + 0.5f);
float ax = x + 1; float av = y + 1;
if (is_maxima<fmaxf>(ax + 0.5f, ay + 0.5f, current))
result[y * width + height] = 1.f;
}
```

However, it gives me a compiler error saying:

```
error: no instance of function template "is_maxima" matches the argument
list
argument types are: (float, float, cudaTextureObject_t)
```

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.

For completeness, `fmaxf`

is defined in the CUDA SDK as:

```
inline float fmaxf(float a, float b)
{
return a > b ? a : b;
}
```