two nearly identical calls, one works one fails.

Question

I have these template functions for use inline on device with cuda

template <class T> __device__ inline T& cmin(T&a,T&b){return (a<b)?(a):(b);};
template <class T> __device__ inline T& cmax(T&a,T&b){return (a>b)?(a):(b);};

In the code I have

cmin(z[i],y[j])-cmax(x[i],z[j])

for int arrays x,y,and z. I get the error:

error: no instance of function template "cmax" matches the argument list

      argument types are: (int, int)

I get the error for cmax but not cmin. If I replace the cmax line with

#define cmax(a,b) ((a)>(b))?(a):(b)

that works just fine, but I don't want #defines, they are problematic. What the heck is going on here?

EDIT: here is the full calling function. times is typedef int.

__global__ void compute_integral_y_isums(times * admit, times * discharge, times * Mx, times * isums, ar_size N){
    // computes the sums for each j
    // blocks on j,
    // threads on i since we will be summing over i.
    // sumation over J should be handled by either a different kernel or on the cpu.
    ar_index tid = threadIdx.x;
    ar_index i = blockIdx.x;                       // summing for patient i 
    ar_index j = threadIdx.x; // across other patients j
    __shared__ times cache[threadsPerBlock];

  times Iy = 0;
    while(j<N){
        // R code:  max(0,min(Mx[i],d3[j,'Discharge.time'])-max(d3[i,'Admission.time'],Mx[j]))
        times imin = cmin(Mx[i],discharge[j]);
        times imax = cmax(admit[i],Mx[j]);
        Iy += cmax(0,imin-imax);
        j += blockDim.x;
    }
    cache[tid] = Iy;

    __syncthreads(); 
    // reduce 
    /***REMOVED***/
}

Answer 1

Try to reverse the definition order please.

template <class T> __device__ inline T& cmax(T&a,T&b){return (a>b)?(a):(b);};
template <class T> __device__ inline T& cmin(T&a,T&b){return (a<b)?(a):(b);};

cmax then cmin. What is the outpout then ?

Answer 2

Your cmax and cmin are taking non-const reference to the elements. Maybe your arrays are declared as const?

Hard to tell, because the example is not complete.

Answer 3

If either x or z is a const array, their element type will be const int, which is not convertible to int&.

Try with:

template<class T> __device__ inline T cmin(const T& a, const T& b)
{
    return (a < b ? a : b);
}

template<class T> __device__ inline T cmax(const T& a, const T& b)
{
    return (a > b ? a : b);
}

If T is always a primitive type like int, you can even pass the parameters by value:

template<class T> __device__ inline T cmin(T a, T b)
{
    return (a < b ? a : b);
}

template<class T> __device__ inline T cmax(T a, T b)
{
    return (a > b ? a : b);
}

EDIT: @aschepler has the right answer.

Answer 4

You should be careful in returning a reference, if your functions also take references as arguments. You might return a reference to a temporary! Like in:

cmin(0,imin-imax);

which is probably ok for int and float, but dangerous for non-PODs.

Answer 5

Iy += cmax(0,imin-imax);

is not legal. You can't bind the literal 0 to an int& reference (but you can to a const int& reference).