diff --git a/modules/gpu/include/opencv2/gpu/device/functional.hpp b/modules/gpu/include/opencv2/gpu/device/functional.hpp index c601cf5273..6e0471e9ac 100644 --- a/modules/gpu/include/opencv2/gpu/device/functional.hpp +++ b/modules/gpu/include/opencv2/gpu/device/functional.hpp @@ -302,18 +302,18 @@ namespace cv { namespace gpu { namespace device template <> struct name : binary_function \ { \ __device__ __forceinline__ type operator()(type lhs, type rhs) const {return op(lhs, rhs);} \ - __device__ __forceinline__ name(const name& other):binary_function(){}\ - __device__ __forceinline__ name():binary_function(){}\ + __device__ __forceinline__ name() {}\ + __device__ __forceinline__ name(const name&) {}\ }; template struct maximum : binary_function { __device__ __forceinline__ T operator()(typename TypeTraits::ParameterType lhs, typename TypeTraits::ParameterType rhs) const { - return lhs < rhs ? rhs : lhs; + return max(lhs, rhs); } - __device__ __forceinline__ maximum(const maximum& other):binary_function(){} - __device__ __forceinline__ maximum():binary_function(){} + __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::ParameterType lhs, typename TypeTraits::ParameterType rhs) const { - return lhs < rhs ? lhs : rhs; + return min(lhs, rhs); } - __device__ __forceinline__ minimum(const minimum& other):binary_function(){} - __device__ __forceinline__ minimum():binary_function(){} + __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 struct abs_func : unary_function + { + __device__ __forceinline__ T operator ()(typename TypeTraits::ParameterType x) const + { + return abs(x); + } + }; + template <> struct abs_func : unary_function + { + __device__ __forceinline__ unsigned char operator ()(unsigned char x) const + { + return x; + } + }; + template <> struct abs_func : unary_function + { + __device__ __forceinline__ signed char operator ()(signed char x) const + { + return ::abs(x); + } + }; + template <> struct abs_func : unary_function + { + __device__ __forceinline__ char operator ()(char x) const + { + return ::abs(x); + } + }; + template <> struct abs_func : unary_function + { + __device__ __forceinline__ unsigned short operator ()(unsigned short x) const + { + return x; + } + }; + template <> struct abs_func : unary_function + { + __device__ __forceinline__ short operator ()(short x) const + { + return ::abs(x); + } + }; + template <> struct abs_func : unary_function + { + __device__ __forceinline__ unsigned int operator ()(unsigned int x) const + { + return x; + } + }; + template <> struct abs_func : unary_function + { + __device__ __forceinline__ int operator ()(int x) const + { + return ::abs(x); + } + }; + template <> struct abs_func : unary_function + { + __device__ __forceinline__ float operator ()(float x) const + { + return ::fabsf(x); + } + }; + template <> struct abs_func : unary_function + { + __device__ __forceinline__ double operator ()(double x) const + { + return ::fabs(x); + } + }; + #define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(name, func) \ template struct name ## _func : unary_function \ { \ @@ -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) diff --git a/modules/gpu/include/opencv2/gpu/device/vec_math.hpp b/modules/gpu/include/opencv2/gpu/device/vec_math.hpp index 0ec790c0b7..1c46dc0c33 100644 --- a/modules/gpu/include/opencv2/gpu/device/vec_math.hpp +++ b/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__ \ No newline at end of file +#endif // __OPENCV_GPU_VECMATH_HPP__ diff --git a/modules/gpu/src/cuda/pyrlk.cu b/modules/gpu/src/cuda/pyrlk.cu index d1a65c210f..811c3b90b5 100644 --- a/modules/gpu/src/cuda/pyrlk.cu +++ b/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 @@ -681,4 +681,4 @@ namespace cv { namespace gpu { namespace device } }}} -#endif /* CUDA_DISABLER */ \ No newline at end of file +#endif /* CUDA_DISABLER */ diff --git a/modules/gpu/src/cuda/surf.cu b/modules/gpu/src/cuda/surf.cu index 8c80559c5b..aebda0ea9e 100644 --- a/modules/gpu/src/cuda/surf.cu +++ b/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 */ \ No newline at end of file +#endif /* CUDA_DISABLER */