два почти идентичных звонка, один работает, другой не работает

Question

У меня есть эти функции шаблона для использования встроенного на устройстве с 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);};

В коде у меня есть

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

для массивов int x, y и z. Я получаю ошибку:

ошибка: ни один экземпляр шаблона функции "cmax" не соответствует списку аргументов

      argument types are: (int, int)

Я получаю ошибку для cmax, но не для cmin. Если я заменю строку cmax на

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

это прекрасно работает, но я не хочу #defines, они проблематичны. Что, черт возьми, здесь происходит?

EDIT: вот полная функция вызова. раз 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

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

не законно. Вы не можете привязать литерал 0 к ссылке int& (но вы можете привязать ссылку const int&).

Answer 2

Если x или z является массивом const, их тип элемента будет const int, который не может быть преобразован в int&.

Попробуйте:

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

Если T всегда является примитивным типом, таким как int, вы даже можете передавать параметры по значению:

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

РЕДАКТИРОВАТЬ: @aschepler имеет правоответить.

Answer 3

Вы должны быть осторожны при возврате ссылки , если ваши функции также принимают ссылки в качестве аргументов.Вы можете вернуть ссылку на временную !Как в:

cmin(0,imin-imax);

, что, вероятно, подходит для int и float, но опасно для не POD.

Answer 4

Чтобы уточнить мой комментарий по первоначальному вопросу: да, нужно быть осторожным, если ona просто проходит через ссылку из параметра.Вот полная иллюстрация:

#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) <b>// return a ref</b>
{
    return a < b ? a : b;
}


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

с последующим

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;
    <b>const Thing &x = refmin(c, b-a); // '&x' will be gone after ';' </b>
    // 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;
}

Вывод:

$ ./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

На некоторых машинах это может даже segfault , наверное.В деструкторе Thing я сбрасываю data на 0, но можно легко представить себе больше.

Итак, когда мы делаем вызов BAD refmin, мы возвращаемсяссылка на временный.Который уничтожен после ;.Итак, когда мы пытаемся вывести &x, оно уже пропало.

Answer 5

Может быть, у вас уже есть активные defines, которые загрязняют ваше пространство имен? Попробуйте переименовать cmin и cmax, или #undef cmin и #undef cmax. Или запустите g++ -E, чтобы увидеть демакрофированный код.

Или добавить :: спецификатор пространства имен:

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

В любом случае, вам нужны только все () в определениях. Nicer:

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

И вам, вероятно, не нужна inline для функции шаблона.

Answer 6

Ваши cmax и cmin принимают неконстантную ссылку на элементы.Может быть, ваши массивы объявлены как const?

Трудно сказать, потому что пример не завершен.

Answer 7

Попробуйте изменить порядок определения, пожалуйста.

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, затем cmin. Какой тогда выход?