CODEWITHSUNDEEP
c++templatescudaoperator-overloading
Vitality
Vitality

Reputation: 21505

Overloading operator= between customized classes int2_, float2_ and double2_ for CUDA application

I'm defining the classes int2_, float2_, and double2_ to deal with complex arithmetics in C++ and CUDA. I want to overload the operator = for mixed assignments of objects of the above classes and of the int, float and double types.

My implementation is the following:

class float2_;
class double2_;

class int2_ {

    public:
        int x;
        int y;

        int2_() : x(), y() {}

        __host__ __device__ inline const int2_& operator=(const int a)          { x = a;            y = 0.;         return *this; }
        __host__ __device__ inline const int2_& operator=(const float a)        { x = (int)a;       y = 0.;         return *this; }
        __host__ __device__ inline const int2_& operator=(const double a)       { x = (int)a;       y = 0.;         return *this; }
        __host__ __device__ inline const int2_& operator=(const int2_ a)        { x = a.x;          y = a.y;        return *this; }
        __host__ __device__ inline const int2_& operator=(const float2_ a);
        __host__ __device__ inline const int2_& operator=(const double2_ a);
};

class float2_ {

    public:
        float x;
        float y;

        float2_() : x(), y() {}

        __host__ __device__ inline const float2_& operator=(const int a)        { x = (float)a;     y = 0.;         return *this; }
        __host__ __device__ inline const float2_& operator=(const float a)      { x = a;            y = 0.;         return *this; }
        __host__ __device__ inline const float2_& operator=(const double a)     { x = (float)a;     y = 0.;         return *this; }
        __host__ __device__ inline const float2_& operator=(const int2_ a)      { x = (float)a.x;   y = (float)a.y; return *this; }
        __host__ __device__ inline const float2_& operator=(const float2_ a)    { x = a.x;          y = a.y;        return *this; }
        __host__ __device__ inline const float2_& operator=(const double2_ a);
};

class double2_ {

    public:
        double x;
        double y;

        double2_() : x(), y() {}

        __host__ __device__ inline const double2_& operator=(const int a)       { x = (double)a;    y = 0.;         return *this; }
        __host__ __device__ inline const double2_& operator=(const float a)     { x = (double)a;    y = 0.;         return *this; }
        __host__ __device__ inline const double2_& operator=(const double a)    { x = a;            y = 0.;         return *this; }
        __host__ __device__ inline const double2_& operator=(const int2_ a)     { x = (double)a.x;  y = (double)a.y;return *this; }
        __host__ __device__ inline const double2_& operator=(const float2_ a)   { x = (double)a.x;  y = (double)a.y;return *this; }
        __host__ __device__ inline const double2_& operator=(const double2_ a)  { x = a.x;          y = a.y;        return *this; }

};

__host__ __device__ inline const int2_& int2_::operator=(const float2_ a)           { x = (int)a.x;             y = (int)a.y;       return *this; }
__host__ __device__ inline const int2_& int2_::operator=(const double2_ a)      { x = (int)a.x;             y = (int)a.y;       return *this; }
__host__ __device__ inline const float2_& float2_::operator=(const double2_ a)  { x = (float)a.x;           y = (float)a.y;     return *this; }

However, I receive a compilation error in the kernel

template <class A, class T1, class T2>
__global__ inline void evaluation_matrix(T1 *data_, const Expr<A,T2> e, int NumElements)
{
    const int i = blockDim.x * blockIdx.x + threadIdx.x;
    if(i < NumElements) data_[i] = e[i];
}

when e is an expression. The error message is

calling a __host__ function("float2_::float2_") from a __global__  
function("evaluation_matrix<BinExpr<const float *, const float2_ *, CudaOpSum, float2_> 
, double2_, float2_> ") is not allowed

In this case, data_ is a double2_ object and e is a float2_ expression.

I have no problem in dealing with any of the int, float, double, int2_, float2_ or double2_ type or class for data_. I even receive no error message when e is an expression of int, float or double type. The only problem arises when e is of class int2_, float2_ or double2_.

Any help? Thank you.

WORKING SOLUTION FOLLOWING ARNE MERTZ'S ANSWER

class float2_;
class double2_;

class int2_ {

    public:
        int x;
        int y;

        __host__ __device__ int2_() : x(), y() {}

        __host__ __device__ inline const int2_& operator=(const int a)          { x = a;            y = 0.;         return *this; }
        __host__ __device__ inline const int2_& operator=(const float a)        { x = (int)a;       y = 0.;         return *this; }
        __host__ __device__ inline const int2_& operator=(const double a)       { x = (int)a;       y = 0.;         return *this; }
        __host__ __device__ inline const int2_& operator=(const int2_ a)        { x = a.x;          y = a.y;        return *this; }
        __host__ __device__ inline const int2_& operator=(const float2_ a);
        __host__ __device__ inline const int2_& operator=(const double2_ a);
};

class float2_ {

    public:
        float x;
        float y;

        __host__ __device__ float2_() : x(), y() {}

        __host__ __device__ inline const float2_& operator=(const int a)        { x = (float)a;     y = 0.;         return *this; }
        __host__ __device__ inline const float2_& operator=(const float a)      { x = a;            y = 0.;         return *this; }
        __host__ __device__ inline const float2_& operator=(const double a)     { x = (float)a;     y = 0.;         return *this; }
        __host__ __device__ inline const float2_& operator=(const int2_ a)      { x = (float)a.x;   y = (float)a.y; return *this; }
        __host__ __device__ inline const float2_& operator=(const float2_ a)    { x = a.x;          y = a.y;        return *this; }
        __host__ __device__ inline const float2_& operator=(const double2_ a);
};

class double2_ {

    public:
        double x;
        double y;

        __host__ __device__ double2_() : x(), y() {}

        __host__ __device__ inline const double2_& operator=(const int a)       { x = (double)a;    y = 0.;         return *this; }
        __host__ __device__ inline const double2_& operator=(const float a)     { x = (double)a;    y = 0.;         return *this; }
        __host__ __device__ inline const double2_& operator=(const double a)    { x = a;            y = 0.;         return *this; }
        __host__ __device__ inline const double2_& operator=(const int2_ a)     { x = (double)a.x;  y = (double)a.y;return *this; }
        __host__ __device__ inline const double2_& operator=(const float2_ a)   { x = (double)a.x;  y = (double)a.y;return *this; }
        __host__ __device__ inline const double2_& operator=(const double2_ a)  { x = a.x;          y = a.y;        return *this; }

};

__host__ __device__ inline const int2_& int2_::operator=(const float2_ a)           { x = (int)a.x;             y = (int)a.y;       return *this; }
__host__ __device__ inline const int2_& int2_::operator=(const double2_ a)      { x = (int)a.x;             y = (int)a.y;       return *this; }
__host__ __device__ inline const float2_& float2_::operator=(const double2_ a)  { x = (float)a.x;           y = (float)a.y;     return *this; }

Upvotes: 0

Views: 931

Answers (1)

Arne Mertz
Arne Mertz

Reputation: 24626

Well, the error says you cannot call __host__ functions (in this case the constructor of your float2_ class) from a __global__ function. At first sight, that has nothing to do with the operators, since they are not mentioned in the error message. But if you have a closer look, there is data_[i] = e[i].

Attention: wild guesses here, since you don't show all of the relevant code:

I guess e[i] gives a reference to a part of the expression, in this case of type float2_. You are assigning that e[i] to a double2_, and the corresponding assignment operator is your double2_::__host__ __device__ inline const double2_& operator=(const float2_ a) which - apart from needlessly and unconventionally returning a const reference, takes that float2_ by value, so the compiler has to copy e[i], by a copy constructor wich seems to be declared with __host__. Apparently from the compiler message, it's not allowed to call __host__ from __global__ functions.

A solution would be to either declare the constructor __global__ or to let the op= take its parameter by (const) reference, so the copy constructor does not need to be called. However, since the operator= itself is declared __host__, you probably get the same error due to that call as well.

I have no idea about cuda and don't know anything more about __host__ and __global__ than that error message told me, but hopefully I could give you a hint what might be wrong with the code.

Upvotes: 2

Related Questions