fix abs_func and minimum/maximum functors

modules/gpu/include/opencv2/gpu/device/functional.hpp

@@ -302,18 +302,18 @@ namespace cv { namespace gpu { namespace device
     template <> struct name<type> : binary_function<type, type, type> \
     { \
         __device__ __forceinline__ type operator()(type lhs, type rhs) const {return op(lhs, rhs);} \
-        __device__ __forceinline__ name(const name& other):binary_function<type, type, type>(){}\
-        __device__ __forceinline__ name():binary_function<type, type, type>(){}\
+        __device__ __forceinline__ name() {}\
+        __device__ __forceinline__ name(const name&) {}\
     };
 
     template <typename T> struct maximum : binary_function<T, T, T>
     {
         __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType lhs, typename TypeTraits<T>::ParameterType rhs) const
         {
-            return lhs < rhs ? rhs : lhs;
+            return max(lhs, rhs);
         }
-        __device__ __forceinline__ maximum(const maximum& other):binary_function<T, T, T>(){}
-        __device__ __forceinline__ maximum():binary_function<T, T, T>(){}
+        __device__ __forceinline__ maximum() {}
+        __device__ __forceinline__ maximum(const maximum&) {}
     };
 
     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, uchar, ::max)
@@ -330,10 +330,10 @@ namespace cv { namespace gpu { namespace device
     {
         __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType lhs, typename TypeTraits<T>::ParameterType rhs) const
         {
-            return lhs < rhs ? lhs : rhs;
+            return min(lhs, rhs);
         }
-        __device__ __forceinline__ minimum(const minimum& other):binary_function<T, T, T>(){}
-        __device__ __forceinline__ minimum():binary_function<T, T, T>(){}
+        __device__ __forceinline__ minimum() {}
+        __device__ __forceinline__ minimum(const minimum&) {}
     };
 
     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, uchar, ::min)
@@ -350,6 +350,78 @@ namespace cv { namespace gpu { namespace device
 
     // Math functions
 ///bound=========================================
+
+    template <typename T> struct abs_func : unary_function<T, T>
+    {
+        __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType x) const
+        {
+            return abs(x);
+        }
+    };
+    template <> struct abs_func<unsigned char> : unary_function<unsigned char, unsigned char>
+    {
+        __device__ __forceinline__ unsigned char operator ()(unsigned char x) const
+        {
+            return x;
+        }
+    };
+    template <> struct abs_func<signed char> : unary_function<signed char, signed char>
+    {
+        __device__ __forceinline__ signed char operator ()(signed char x) const
+        {
+            return ::abs(x);
+        }
+    };
+    template <> struct abs_func<char> : unary_function<char, char>
+    {
+        __device__ __forceinline__ char operator ()(char x) const
+        {
+            return ::abs(x);
+        }
+    };
+    template <> struct abs_func<unsigned short> : unary_function<unsigned short, unsigned short>
+    {
+        __device__ __forceinline__ unsigned short operator ()(unsigned short x) const
+        {
+            return x;
+        }
+    };
+    template <> struct abs_func<short> : unary_function<short, short>
+    {
+        __device__ __forceinline__ short operator ()(short x) const
+        {
+            return ::abs(x);
+        }
+    };
+    template <> struct abs_func<unsigned int> : unary_function<unsigned int, unsigned int>
+    {
+        __device__ __forceinline__ unsigned int operator ()(unsigned int x) const
+        {
+            return x;
+        }
+    };
+    template <> struct abs_func<int> : unary_function<int, int>
+    {
+        __device__ __forceinline__ int operator ()(int x) const
+        {
+            return ::abs(x);
+        }
+    };
+    template <> struct abs_func<float> : unary_function<float, float>
+    {
+        __device__ __forceinline__ float operator ()(float x) const
+        {
+            return ::fabsf(x);
+        }
+    };
+    template <> struct abs_func<double> : unary_function<double, double>
+    {
+        __device__ __forceinline__ double operator ()(double x) const
+        {
+            return ::fabs(x);
+        }
+    };
+
 #define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(name, func) \
     template <typename T> struct name ## _func : unary_function<T, float> \
     { \
@@ -382,7 +454,6 @@ namespace cv { namespace gpu { namespace device
         } \
     };
 
-    OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(fabs, ::fabs)
     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sqrt, ::sqrt)
     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp, ::exp)
     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp2, ::exp2)

modules/gpu/include/opencv2/gpu/device/vec_math.hpp

@@ -280,7 +280,7 @@ namespace cv { namespace gpu { namespace device
     OPENCV_GPU_IMPLEMENT_VEC_UNOP (type, operator ! , logical_not) \
     OPENCV_GPU_IMPLEMENT_VEC_BINOP(type, max, maximum) \
     OPENCV_GPU_IMPLEMENT_VEC_BINOP(type, min, minimum) \
-    OPENCV_GPU_IMPLEMENT_VEC_UNOP(type, fabs, fabs_func) \
+    OPENCV_GPU_IMPLEMENT_VEC_UNOP(type, abs, abs_func) \
     OPENCV_GPU_IMPLEMENT_VEC_UNOP(type, sqrt, sqrt_func) \
     OPENCV_GPU_IMPLEMENT_VEC_UNOP(type, exp, exp_func) \
     OPENCV_GPU_IMPLEMENT_VEC_UNOP(type, exp2, exp2_func) \
@@ -327,4 +327,4 @@ namespace cv { namespace gpu { namespace device
     #undef OPENCV_GPU_IMPLEMENT_VEC_INT_OP
 }}} // namespace cv { namespace gpu { namespace device
 
-#endif // __OPENCV_GPU_VECMATH_HPP__
+#endif // __OPENCV_GPU_VECMATH_HPP__

modules/gpu/src/cuda/pyrlk.cu

@@ -267,7 +267,7 @@ namespace cv { namespace gpu { namespace device
         }
         __device__ __forceinline__ float4 abs_(const float4& a)
         {
-            return fabs(a);
+            return abs(a);
         }
 
         template <int cn, int PATCH_X, int PATCH_Y, bool calcErr>
@@ -681,4 +681,4 @@ namespace cv { namespace gpu { namespace device
     }
 }}}
 
-#endif /* CUDA_DISABLER */
+#endif /* CUDA_DISABLER */

modules/gpu/src/cuda/surf.cu

@@ -638,7 +638,7 @@ namespace cv { namespace gpu { namespace device
                 kp_dir *= 180.0f / CV_PI_F;
 
                 kp_dir = 360.0f - kp_dir;
-                if (abs(kp_dir - 360.f) < FLT_EPSILON)
+                if (::fabsf(kp_dir - 360.f) < FLT_EPSILON)
                     kp_dir = 0.f;
 
                 featureDir[blockIdx.x] = kp_dir;
@@ -1003,4 +1003,4 @@ namespace cv { namespace gpu { namespace device
 }}} // namespace cv { namespace gpu { namespace device
 
 
-#endif /* CUDA_DISABLER */
+#endif /* CUDA_DISABLER */