Facing errors with cuda operator overloading

I have cuda kernel that does something like below

__device int func(const float4 vec)
 {
   float4 test = -vec;
 }

This gave me compilation error about no operator - takes argument const float4 so I overloaded operator - like below

inline __device__ float4 operator-(const float4 a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}

That fixed the error but now I am getting some more errors.

I have code like below

 __device int func()
 {
   float4 test = ...;
   test = -test;
 }

Now I am getting following errors

 error: more than one operator "-" matches these operands:
            function "operator-(float4 &)"
            function "operator-(float4)"
            operand types are: - float4

any one know how I can fix this ?

Thanks

After fixing some typos, I don’t have any trouble with what you have shown:

$ cat t30.cu
inline __device__ float4 operator-(const float4 a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}

__device__ void func1(const float4 vec)
 {
   float4 test = -vec;
 }

 __device__ void func2(const float4 vec)
 {
   float4 test = vec;
   test = -test;
 }
$ nvcc -c t30.cu
$
1 Like

My guess would be that you actually have something like this:

$ cat t30.cu
inline __device__ float4 operator-(const float4 a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}
inline __device__ float4 operator-(const float4 &a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}

__device__ void func1(const float4 vec)
 {
   float4 test = -vec;
 }

 __device__ void func2(const float4 vec)
 {
   float4 test = vec;
   test = -test;
 }
$ nvcc -c t30.cu
t30.cu(6): error: more than one operator "-" matches these operands:
            function "operator-(float4)"
            function "operator-(const float4 &)"
            operand types are: - const float4

t30.cu(12): error: more than one operator "-" matches these operands:
            function "operator-(float4)"
            function "operator-(const float4 &)"
            operand types are: - float4

2 errors detected in the compilation of "t30.cu".
$

In which case I would say “don’t do that.” Get rid of the operator overload that passes its argument by value, and just use the pass-by-reference version.

By the way I doubt this problem has anything to do with CUDA:

$ cat t30.cpp
#define __device__
#include "/usr/local/cuda/include/vector_types.h"
#include "/usr/local/cuda/include/vector_functions.h"
inline __device__ float4 operator-(const float4 a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}
inline __device__ float4 operator-(const float4 &a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}

__device__ void func1(const float4 vec)
 {
   float4 test = -vec;
 }

 __device__ void func2(const float4 vec)
 {
   float4 test = vec;
   test = -test;
 }
$ g++ -c t30.cpp
t30.cpp: In function ‘void func1(float4)’:
t30.cpp:9:18: error: ambiguous overload for ‘operator-’ (operand type is ‘const float4’)
    float4 test = -vec;
                  ^~~~
t30.cpp:4:26: note: candidate: float4 operator-(float4)
 inline __device__ float4 operator-(const float4 a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}
                          ^~~~~~~~
t30.cpp:5:26: note: candidate: float4 operator-(const float4&)
 inline __device__ float4 operator-(const float4 &a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}
                          ^~~~~~~~
t30.cpp: In function ‘void func2(float4)’:
t30.cpp:15:11: error: ambiguous overload for ‘operator-’ (operand type is ‘float4’)
    test = -test;
           ^~~~~
t30.cpp:4:26: note: candidate: float4 operator-(float4)
 inline __device__ float4 operator-(const float4 a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}
                          ^~~~~~~~
t30.cpp:5:26: note: candidate: float4 operator-(const float4&)
 inline __device__ float4 operator-(const float4 &a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}
                          ^~~~~~~~
$

I did not have

inline __device__ float4 operator-(const float4 a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}
inline __device__ float4 operator-(const float4 &a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}

but I had only one as I showed in code snippet.

inline __device__ float4 operator-(const float4 a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}

But I replaced above function with

inline __device__ float4 operator-(const float4& a) {return make_float4(-a.x,-a.y,-a.z,-a.w);}

That fixed my issue.

Thanks.

The other one was around, somewhere. Your error output makes that evident, and I’ve already demonstrated that your code snippet, by itself, does not generate the error you are showing. If you still disagree, my suggestion would be to provide a complete example that I can compile and see the error output you are suggesting.