cuda device function and templates

Question

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;
}

here we go again, lacking an MCVE. You've got something that *already* reproduces the error. Why not provide a complete code? If you think it's because you're asking a "conceptual" question, you're mistaken, because every debugging question can be mapped into a conceptual question for the purposes of avoiding the MCVE requirement. — Robert Crovella, May 26 '15 at 16:37
the error message seems to be complaining about the argument list. Is `current` a valid cuda texture object type? Also, your template parameter is Op with a capital O, and your code is using op with a small o. Doesn't look right. — Robert Crovella, May 26 '15 at 16:42
The small 'o' was a type from me. I corrected it. The texture is valid. I am getting the error at the compilation stage anyway. — Luca, May 26 '15 at 16:44
Ok, I got it. The signature should have been: template __device__ bool is_maxima(float ax, float ay, cudaTextureObject_t current) and the usage if (op(cv, pv) != cv) return false; — Luca, May 26 '15 at 16:49
According to my testing it should be `template` not `template` — Robert Crovella, May 26 '15 at 17:01

score 3 · Accepted Answer · answered May 26 '15 at 17:05

As you've already pointed out, this is not correct:

template<typename Op>

I believe it should be:

template<Op op>

The following code seems to work correctly for me (removing the texturing since it's extraneous to the issue):

$ cat t755.cu
#include <stdio.h>
#include <math.h>

typedef float(*Op)(float, float);


__device__ float mymax(float a, float b){

  return (a>b)?a:b;
}

template <Op op>
__device__ float test(float d1, float d2){
  return op(d1, d2);
}

__global__ void kernel(){

  float a1 = 1.0f;
  float a2 = 2.0f;
  printf("%f\n", test<mymax>(a1, a2));
}

int main(){

  kernel<<<1,1>>>();
  cudaDeviceSynchronize();
}
$ nvcc -arch=sm_35 -std=c++11 t755.cu -o t755
$ ./t755
2.000000
$

yeah, thanks. I figured that out as well after many trial and error. — Luca, May 26 '15 at 17:45
also helpful to check https://stackoverflow.com/questions/34879789/thrust-transform-throws-error-bulk-kernel-by-value-an-illegal-memory-access-w — Izana, Jan 19 '23 at 00:29

score 1 · Answer 2 · edited May 23 '17 at 11:58

You could define is_maxima like this in order to avoid the typedef:

template<typename Op>
__device__ bool is_maxima(float ax, float ay, cudaTextureObject_t current, Op fun)
{
    float cv = tex2D<float>(current, ax, ay);    
    float pv = tex2D<float>(current, ax - 1.f, ay);
    if (fun(cv, pv) != cv) return false;
    return true;
}

And then call it using is_maxima(ax + 0.5f, ay + 0.5f, current, fmaxf).

See this answer for more details: https://stackoverflow.com/a/1174193/678093

cuda device function and templates

2 Answers2