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

To clarify my comment on the initial question: Yes, one should be careful if ona just passes through a reference from a parameter. Here is a complete illustration:

#include <iostream>

struct Thing {
    int data;
    Thing() { data = 42; }
    Thing(int val) { data = val; }
    Thing(const Thing& oth) { data = oth.data; }
    Thing& operator=(const Thing& oth)
      { if(this!=&oth) this->data = oth.data; return *this; }
    ~Thing() { data = 0; } // clear, destroy, release...
};

bool operator<(const Thing &a, const  Thing &b) { return a.data 
const T& refmin(const T &a, const T &b) // return a ref
{
    return a < b ? a : b;
}


template
const T copymin(const T &a, const T &b) // return a copy
{
    return a < b ? a : b;
}

followed by

int main(int argc, const char* [])
{
    Thing a(11);
    Thing b(88);

    std::cerr << "Simple operation:" << std::endl;
    const Thing c = a + b;
    std::cerr << "  c:" << c.data << "  should be 99" << std::endl;

    std::cerr << "Working on temp expression (BAD):" << std::endl;
    const Thing &x = refmin(c, b-a); // '&x' will be gone after ';' 
    // the next line might crash:
    std::cerr << "  x:" << x.data << "  should be 77" << std::endl;

    std::cerr << "Working explicit side variable (OK):" << std::endl;
    const Thing &d = b-a;
    const Thing &y = refmin(c, d);  // '&y' is now same as '&d'
    std::cerr << "  d:" << d.data << "  should be 77" << std::endl;
    std::cerr << "  y:" << y.data << "  should be 77" << std::endl;

    std::cerr << "Working on a copy (OK):" << std::endl;
    const Thing &z = copymin(c, b-a);
    std::cerr <<  z:" << z.data << "  should be 77" << std::endl;

    return 0;
}

The output is:

$ ./minreftest
Simple operation:
  c:99  should be 99
Working on temp expression (BAD):
  x:0  should be 77
Working explicit side variable (OK):
  d:77  should be 77
  y:77  should be 77
Working on a copy (OK):
  z:77  should be 77

On some machines it might even segfault, I guess. In Things destructor I reset data to 0, but one could easily imagine more there.

So, when we do the BAD refmin call, we return a reference to a temporary. Which is destroyed after the ;. So, when we try to output &x, it's already gone.

Answer 2

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 3

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).

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

Maybe you already have some active defines which pollute you namespace?Try renaming cmin and cmax, or #undef cmin and #undef cmax. Or run g++ -E to see the de-macrofied code.

Or add :: namespace specifier:

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

Anyway, you only need all the () in defines. Nicer:

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

And you probably do not need the inline for a template function neither.

Answer 6

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 7

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 ?