1 /*M/////////////////////////////////////////////////////////////////////////////////////// 2 // 3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. 4 // 5 // By downloading, copying, installing or using the software you agree to this license. 6 // If you do not agree to this license, do not download, install, 7 // copy or use the software. 8 // 9 // 10 // License Agreement 11 // For Open Source Computer Vision Library 12 // 13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved. 14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved. 15 // Copyright (C) 2013, OpenCV Foundation, all rights reserved. 16 // Third party copyrights are property of their respective owners. 17 // 18 // Redistribution and use in source and binary forms, with or without modification, 19 // are permitted provided that the following conditions are met: 20 // 21 // * Redistribution's of source code must retain the above copyright notice, 22 // this list of conditions and the following disclaimer. 23 // 24 // * Redistribution's in binary form must reproduce the above copyright notice, 25 // this list of conditions and the following disclaimer in the documentation 26 // and/or other materials provided with the distribution. 27 // 28 // * The name of the copyright holders may not be used to endorse or promote products 29 // derived from this software without specific prior written permission. 30 // 31 // This software is provided by the copyright holders and contributors "as is" and 32 // any express or implied warranties, including, but not limited to, the implied 33 // warranties of merchantability and fitness for a particular purpose are disclaimed. 34 // In no event shall the Intel Corporation or contributors be liable for any direct, 35 // indirect, incidental, special, exemplary, or consequential damages 36 // (including, but not limited to, procurement of substitute goods or services; 37 // loss of use, data, or profits; or business interruption) however caused 38 // and on any theory of liability, whether in contract, strict liability, 39 // or tort (including negligence or otherwise) arising in any way out of 40 // the use of this software, even if advised of the possibility of such damage. 41 // 42 //M*/ 43 44 #pragma once 45 46 #ifndef __OPENCV_CUDEV_UTIL_VEC_MATH_HPP__ 47 #define __OPENCV_CUDEV_UTIL_VEC_MATH_HPP__ 48 49 #include "vec_traits.hpp" 50 #include "saturate_cast.hpp" 51 52 namespace cv { namespace cudev { 53 54 //! @addtogroup cudev 55 //! @{ 56 57 // saturate_cast 58 59 namespace vec_math_detail 60 { 61 template <int cn, typename VecD> struct SatCastHelper; 62 63 template <typename VecD> struct SatCastHelper<1, VecD> 64 { 65 template <typename VecS> __device__ __forceinline__ static VecD cast(const VecS& v) 66 { 67 typedef typename VecTraits<VecD>::elem_type D; 68 return VecTraits<VecD>::make(saturate_cast<D>(v.x)); 69 } 70 }; 71 72 template <typename VecD> struct SatCastHelper<2, VecD> 73 { 74 template <typename VecS> __device__ __forceinline__ static VecD cast(const VecS& v) 75 { 76 typedef typename VecTraits<VecD>::elem_type D; 77 return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y)); 78 } 79 }; 80 81 template <typename VecD> struct SatCastHelper<3, VecD> 82 { 83 template <typename VecS> __device__ __forceinline__ static VecD cast(const VecS& v) 84 { 85 typedef typename VecTraits<VecD>::elem_type D; 86 return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y), saturate_cast<D>(v.z)); 87 } 88 }; 89 90 template <typename VecD> struct SatCastHelper<4, VecD> 91 { 92 template <typename VecS> __device__ __forceinline__ static VecD cast(const VecS& v) 93 { 94 typedef typename VecTraits<VecD>::elem_type D; 95 return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y), saturate_cast<D>(v.z), saturate_cast<D>(v.w)); 96 } 97 }; 98 } 99 100 template<typename T> __device__ __forceinline__ T saturate_cast(const uchar1& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 101 template<typename T> __device__ __forceinline__ T saturate_cast(const char1& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 102 template<typename T> __device__ __forceinline__ T saturate_cast(const ushort1& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 103 template<typename T> __device__ __forceinline__ T saturate_cast(const short1& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 104 template<typename T> __device__ __forceinline__ T saturate_cast(const uint1& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 105 template<typename T> __device__ __forceinline__ T saturate_cast(const int1& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 106 template<typename T> __device__ __forceinline__ T saturate_cast(const float1& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 107 template<typename T> __device__ __forceinline__ T saturate_cast(const double1& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 108 109 template<typename T> __device__ __forceinline__ T saturate_cast(const uchar2& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 110 template<typename T> __device__ __forceinline__ T saturate_cast(const char2& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 111 template<typename T> __device__ __forceinline__ T saturate_cast(const ushort2& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 112 template<typename T> __device__ __forceinline__ T saturate_cast(const short2& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 113 template<typename T> __device__ __forceinline__ T saturate_cast(const uint2& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 114 template<typename T> __device__ __forceinline__ T saturate_cast(const int2& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 115 template<typename T> __device__ __forceinline__ T saturate_cast(const float2& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 116 template<typename T> __device__ __forceinline__ T saturate_cast(const double2& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 117 118 template<typename T> __device__ __forceinline__ T saturate_cast(const uchar3& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 119 template<typename T> __device__ __forceinline__ T saturate_cast(const char3& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 120 template<typename T> __device__ __forceinline__ T saturate_cast(const ushort3& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 121 template<typename T> __device__ __forceinline__ T saturate_cast(const short3& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 122 template<typename T> __device__ __forceinline__ T saturate_cast(const uint3& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 123 template<typename T> __device__ __forceinline__ T saturate_cast(const int3& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 124 template<typename T> __device__ __forceinline__ T saturate_cast(const float3& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 125 template<typename T> __device__ __forceinline__ T saturate_cast(const double3& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 126 127 template<typename T> __device__ __forceinline__ T saturate_cast(const uchar4& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 128 template<typename T> __device__ __forceinline__ T saturate_cast(const char4& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 129 template<typename T> __device__ __forceinline__ T saturate_cast(const ushort4& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 130 template<typename T> __device__ __forceinline__ T saturate_cast(const short4& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 131 template<typename T> __device__ __forceinline__ T saturate_cast(const uint4& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 132 template<typename T> __device__ __forceinline__ T saturate_cast(const int4& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 133 template<typename T> __device__ __forceinline__ T saturate_cast(const float4& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 134 template<typename T> __device__ __forceinline__ T saturate_cast(const double4& v) { return vec_math_detail::SatCastHelper<VecTraits<T>::cn, T>::cast(v); } 135 136 // unary operators 137 138 #define CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(op, input_type, output_type) \ 139 __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a) \ 140 { \ 141 return VecTraits<output_type ## 1>::make(op (a.x)); \ 142 } \ 143 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a) \ 144 { \ 145 return VecTraits<output_type ## 2>::make(op (a.x), op (a.y)); \ 146 } \ 147 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a) \ 148 { \ 149 return VecTraits<output_type ## 3>::make(op (a.x), op (a.y), op (a.z)); \ 150 } \ 151 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a) \ 152 { \ 153 return VecTraits<output_type ## 4>::make(op (a.x), op (a.y), op (a.z), op (a.w)); \ 154 } 155 156 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, char, char) 157 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, short, short) 158 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, int, int) 159 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, float, float) 160 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, double, double) 161 162 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, uchar, uchar) 163 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, char, uchar) 164 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, ushort, uchar) 165 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, short, uchar) 166 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, int, uchar) 167 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, uint, uchar) 168 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, float, uchar) 169 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, double, uchar) 170 171 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, uchar, uchar) 172 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, char, char) 173 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, ushort, ushort) 174 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, short, short) 175 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, int, int) 176 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, uint, uint) 177 178 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_OP 179 180 // unary functions 181 182 #define CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(func_name, func, input_type, output_type) \ 183 __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a) \ 184 { \ 185 return VecTraits<output_type ## 1>::make(func (a.x)); \ 186 } \ 187 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a) \ 188 { \ 189 return VecTraits<output_type ## 2>::make(func (a.x), func (a.y)); \ 190 } \ 191 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a) \ 192 { \ 193 return VecTraits<output_type ## 3>::make(func (a.x), func (a.y), func (a.z)); \ 194 } \ 195 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a) \ 196 { \ 197 return VecTraits<output_type ## 4>::make(func (a.x), func (a.y), func (a.z), func (a.w)); \ 198 } 199 200 namespace vec_math_detail 201 { 202 __device__ __forceinline__ schar abs_(schar val) 203 { 204 return (schar) ::abs((int) val); 205 } 206 207 __device__ __forceinline__ short abs_(short val) 208 { 209 return (short) ::abs((int) val); 210 } 211 } 212 213 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, uchar, uchar) 214 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, vec_math_detail::abs_, char, char) 215 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, ushort, ushort) 216 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, vec_math_detail::abs_, short, short) 217 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::abs, int, int) 218 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, uint, uint) 219 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::fabsf, float, float) 220 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::fabs, double, double) 221 222 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, uchar, float) 223 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, char, float) 224 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, ushort, float) 225 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, short, float) 226 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, int, float) 227 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, uint, float) 228 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, float, float) 229 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrt, double, double) 230 231 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, uchar, float) 232 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, char, float) 233 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, ushort, float) 234 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, short, float) 235 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, int, float) 236 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, uint, float) 237 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, float, float) 238 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::exp, double, double) 239 240 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, uchar, float) 241 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, char, float) 242 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, ushort, float) 243 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, short, float) 244 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, int, float) 245 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, uint, float) 246 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, float, float) 247 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2, double, double) 248 249 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, uchar, float) 250 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, char, float) 251 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, ushort, float) 252 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, short, float) 253 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, int, float) 254 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, uint, float) 255 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, float, float) 256 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10, double, double) 257 258 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, uchar, float) 259 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, char, float) 260 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, ushort, float) 261 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, short, float) 262 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, int, float) 263 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, uint, float) 264 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, float, float) 265 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::log, double, double) 266 267 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, uchar, float) 268 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, char, float) 269 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, ushort, float) 270 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, short, float) 271 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, int, float) 272 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, uint, float) 273 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, float, float) 274 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2, double, double) 275 276 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, uchar, float) 277 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, char, float) 278 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, ushort, float) 279 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, short, float) 280 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, int, float) 281 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, uint, float) 282 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, float, float) 283 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10, double, double) 284 285 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, uchar, float) 286 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, char, float) 287 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, ushort, float) 288 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, short, float) 289 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, int, float) 290 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, uint, float) 291 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, float, float) 292 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sin, double, double) 293 294 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, uchar, float) 295 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, char, float) 296 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, ushort, float) 297 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, short, float) 298 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, int, float) 299 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, uint, float) 300 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, float, float) 301 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cos, double, double) 302 303 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, uchar, float) 304 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, char, float) 305 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, ushort, float) 306 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, short, float) 307 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, int, float) 308 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, uint, float) 309 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, float, float) 310 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tan, double, double) 311 312 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, uchar, float) 313 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, char, float) 314 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, ushort, float) 315 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, short, float) 316 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, int, float) 317 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, uint, float) 318 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, float, float) 319 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asin, double, double) 320 321 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, uchar, float) 322 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, char, float) 323 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, ushort, float) 324 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, short, float) 325 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, int, float) 326 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, uint, float) 327 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, float, float) 328 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acos, double, double) 329 330 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, uchar, float) 331 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, char, float) 332 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, ushort, float) 333 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, short, float) 334 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, int, float) 335 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, uint, float) 336 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, float, float) 337 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atan, double, double) 338 339 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, uchar, float) 340 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, char, float) 341 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, ushort, float) 342 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, short, float) 343 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, int, float) 344 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, uint, float) 345 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, float, float) 346 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinh, double, double) 347 348 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, uchar, float) 349 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, char, float) 350 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, ushort, float) 351 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, short, float) 352 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, int, float) 353 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, uint, float) 354 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, float, float) 355 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::cosh, double, double) 356 357 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, uchar, float) 358 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, char, float) 359 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, ushort, float) 360 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, short, float) 361 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, int, float) 362 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, uint, float) 363 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, float, float) 364 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanh, double, double) 365 366 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, uchar, float) 367 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, char, float) 368 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, ushort, float) 369 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, short, float) 370 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, int, float) 371 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, uint, float) 372 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, float, float) 373 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinh, double, double) 374 375 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, uchar, float) 376 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, char, float) 377 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, ushort, float) 378 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, short, float) 379 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, int, float) 380 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, uint, float) 381 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, float, float) 382 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acosh, double, double) 383 384 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, uchar, float) 385 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, char, float) 386 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, ushort, float) 387 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, short, float) 388 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, int, float) 389 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, uint, float) 390 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, float, float) 391 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanh, double, double) 392 393 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC 394 395 // binary operators (vec & vec) 396 397 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(op, input_type, output_type) \ 398 __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, const input_type ## 1 & b) \ 399 { \ 400 return VecTraits<output_type ## 1>::make(a.x op b.x); \ 401 } \ 402 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, const input_type ## 2 & b) \ 403 { \ 404 return VecTraits<output_type ## 2>::make(a.x op b.x, a.y op b.y); \ 405 } \ 406 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, const input_type ## 3 & b) \ 407 { \ 408 return VecTraits<output_type ## 3>::make(a.x op b.x, a.y op b.y, a.z op b.z); \ 409 } \ 410 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, const input_type ## 4 & b) \ 411 { \ 412 return VecTraits<output_type ## 4>::make(a.x op b.x, a.y op b.y, a.z op b.z, a.w op b.w); \ 413 } 414 415 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, uchar, int) 416 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, char, int) 417 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, ushort, int) 418 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, short, int) 419 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, int, int) 420 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, uint, uint) 421 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, float, float) 422 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, double, double) 423 424 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uchar, int) 425 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, char, int) 426 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, ushort, int) 427 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, short, int) 428 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, int, int) 429 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uint, uint) 430 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, float, float) 431 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, double, double) 432 433 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, uchar, int) 434 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, char, int) 435 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, ushort, int) 436 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, short, int) 437 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, int, int) 438 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, uint, uint) 439 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, float, float) 440 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, double, double) 441 442 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, uchar, int) 443 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, char, int) 444 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, ushort, int) 445 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, short, int) 446 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, int, int) 447 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, uint, uint) 448 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, float, float) 449 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, double, double) 450 451 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uchar, uchar) 452 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, char, uchar) 453 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, ushort, uchar) 454 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, short, uchar) 455 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, int, uchar) 456 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uint, uchar) 457 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, float, uchar) 458 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, double, uchar) 459 460 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uchar, uchar) 461 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, char, uchar) 462 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, ushort, uchar) 463 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, short, uchar) 464 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, int, uchar) 465 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uint, uchar) 466 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, float, uchar) 467 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, double, uchar) 468 469 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, uchar, uchar) 470 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, char, uchar) 471 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, ushort, uchar) 472 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, short, uchar) 473 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, int, uchar) 474 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, uint, uchar) 475 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, float, uchar) 476 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, double, uchar) 477 478 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, uchar, uchar) 479 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, char, uchar) 480 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, ushort, uchar) 481 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, short, uchar) 482 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, int, uchar) 483 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, uint, uchar) 484 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, float, uchar) 485 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, double, uchar) 486 487 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, uchar, uchar) 488 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, char, uchar) 489 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, ushort, uchar) 490 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, short, uchar) 491 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, int, uchar) 492 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, uint, uchar) 493 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, float, uchar) 494 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, double, uchar) 495 496 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, uchar, uchar) 497 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, char, uchar) 498 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, ushort, uchar) 499 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, short, uchar) 500 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, int, uchar) 501 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, uint, uchar) 502 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, float, uchar) 503 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, double, uchar) 504 505 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, uchar, uchar) 506 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, char, uchar) 507 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, ushort, uchar) 508 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, short, uchar) 509 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, int, uchar) 510 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, uint, uchar) 511 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, float, uchar) 512 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, double, uchar) 513 514 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, uchar, uchar) 515 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, char, uchar) 516 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, ushort, uchar) 517 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, short, uchar) 518 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, int, uchar) 519 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, uint, uchar) 520 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, float, uchar) 521 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, double, uchar) 522 523 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, uchar, uchar) 524 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, char, char) 525 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, ushort, ushort) 526 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, short, short) 527 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, int, int) 528 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, uint, uint) 529 530 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, uchar, uchar) 531 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, char, char) 532 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, ushort, ushort) 533 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, short, short) 534 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, int, int) 535 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, uint, uint) 536 537 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uchar, uchar) 538 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, char, char) 539 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, ushort, ushort) 540 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, short, short) 541 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, int, int) 542 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uint, uint) 543 544 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_OP 545 546 // binary operators (vec & scalar) 547 548 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(op, input_type, scalar_type, output_type) \ 549 __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, scalar_type s) \ 550 { \ 551 return VecTraits<output_type ## 1>::make(a.x op s); \ 552 } \ 553 __device__ __forceinline__ output_type ## 1 operator op(scalar_type s, const input_type ## 1 & b) \ 554 { \ 555 return VecTraits<output_type ## 1>::make(s op b.x); \ 556 } \ 557 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, scalar_type s) \ 558 { \ 559 return VecTraits<output_type ## 2>::make(a.x op s, a.y op s); \ 560 } \ 561 __device__ __forceinline__ output_type ## 2 operator op(scalar_type s, const input_type ## 2 & b) \ 562 { \ 563 return VecTraits<output_type ## 2>::make(s op b.x, s op b.y); \ 564 } \ 565 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, scalar_type s) \ 566 { \ 567 return VecTraits<output_type ## 3>::make(a.x op s, a.y op s, a.z op s); \ 568 } \ 569 __device__ __forceinline__ output_type ## 3 operator op(scalar_type s, const input_type ## 3 & b) \ 570 { \ 571 return VecTraits<output_type ## 3>::make(s op b.x, s op b.y, s op b.z); \ 572 } \ 573 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, scalar_type s) \ 574 { \ 575 return VecTraits<output_type ## 4>::make(a.x op s, a.y op s, a.z op s, a.w op s); \ 576 } \ 577 __device__ __forceinline__ output_type ## 4 operator op(scalar_type s, const input_type ## 4 & b) \ 578 { \ 579 return VecTraits<output_type ## 4>::make(s op b.x, s op b.y, s op b.z, s op b.w); \ 580 } 581 582 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, int, int) 583 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, float, float) 584 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, double, double) 585 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, int, int) 586 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, float, float) 587 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, double, double) 588 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, int, int) 589 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, float, float) 590 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, double, double) 591 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, int, int) 592 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, float, float) 593 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, double, double) 594 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, int, int) 595 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, float, float) 596 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, double, double) 597 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, uint, uint) 598 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, float, float) 599 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, double, double) 600 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, float, float, float) 601 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, float, double, double) 602 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, double, double, double) 603 604 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, int, int) 605 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, float, float) 606 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, double, double) 607 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, int, int) 608 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, float, float) 609 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, double, double) 610 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, int, int) 611 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, float, float) 612 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, double, double) 613 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, int, int) 614 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, float, float) 615 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, double, double) 616 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, int, int) 617 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, float, float) 618 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, double, double) 619 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, uint, uint) 620 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, float, float) 621 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, double, double) 622 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, float, float, float) 623 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, float, double, double) 624 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, double, double, double) 625 626 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, int, int) 627 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, float, float) 628 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, double, double) 629 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, int, int) 630 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, float, float) 631 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, double, double) 632 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, int, int) 633 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, float, float) 634 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, double, double) 635 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, int, int) 636 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, float, float) 637 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, double, double) 638 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, int, int) 639 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, float, float) 640 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, double, double) 641 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, uint, uint) 642 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, float, float) 643 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, double, double) 644 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, float, float, float) 645 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, float, double, double) 646 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, double, double, double) 647 648 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, int, int) 649 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, float, float) 650 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, double, double) 651 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, int, int) 652 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, float, float) 653 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, double, double) 654 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, int, int) 655 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, float, float) 656 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, double, double) 657 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, int, int) 658 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, float, float) 659 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, double, double) 660 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, int, int) 661 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, float, float) 662 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, double, double) 663 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, uint, uint) 664 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, float, float) 665 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, double, double) 666 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, float, float, float) 667 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, float, double, double) 668 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, double, double, double) 669 670 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uchar, uchar, uchar) 671 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, char, char, uchar) 672 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, ushort, ushort, uchar) 673 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, short, short, uchar) 674 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, int, int, uchar) 675 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uint, uint, uchar) 676 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, float, float, uchar) 677 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, double, double, uchar) 678 679 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uchar, uchar, uchar) 680 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, char, char, uchar) 681 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, ushort, ushort, uchar) 682 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, short, short, uchar) 683 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, int, int, uchar) 684 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uint, uint, uchar) 685 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, float, float, uchar) 686 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, double, double, uchar) 687 688 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, uchar, uchar, uchar) 689 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, char, char, uchar) 690 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, ushort, ushort, uchar) 691 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, short, short, uchar) 692 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, int, int, uchar) 693 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, uint, uint, uchar) 694 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, float, float, uchar) 695 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, double, double, uchar) 696 697 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, uchar, uchar, uchar) 698 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, char, char, uchar) 699 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, ushort, ushort, uchar) 700 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, short, short, uchar) 701 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, int, int, uchar) 702 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, uint, uint, uchar) 703 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, float, float, uchar) 704 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, double, double, uchar) 705 706 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, uchar, uchar, uchar) 707 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, char, char, uchar) 708 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, ushort, ushort, uchar) 709 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, short, short, uchar) 710 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, int, int, uchar) 711 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, uint, uint, uchar) 712 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, float, float, uchar) 713 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, double, double, uchar) 714 715 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, uchar, uchar, uchar) 716 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, char, char, uchar) 717 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, ushort, ushort, uchar) 718 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, short, short, uchar) 719 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, int, int, uchar) 720 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, uint, uint, uchar) 721 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, float, float, uchar) 722 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, double, double, uchar) 723 724 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, uchar, uchar, uchar) 725 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, char, char, uchar) 726 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, ushort, ushort, uchar) 727 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, short, short, uchar) 728 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, int, int, uchar) 729 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, uint, uint, uchar) 730 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, float, float, uchar) 731 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, double, double, uchar) 732 733 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, uchar, uchar, uchar) 734 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, char, char, uchar) 735 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, ushort, ushort, uchar) 736 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, short, short, uchar) 737 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, int, int, uchar) 738 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, uint, uint, uchar) 739 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, float, float, uchar) 740 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, double, double, uchar) 741 742 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, uchar, uchar, uchar) 743 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, char, char, char) 744 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, ushort, ushort, ushort) 745 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, short, short, short) 746 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, int, int, int) 747 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, uint, uint, uint) 748 749 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, uchar, uchar, uchar) 750 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, char, char, char) 751 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, ushort, ushort, ushort) 752 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, short, short, short) 753 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, int, int, int) 754 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, uint, uint, uint) 755 756 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uchar, uchar, uchar) 757 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, char, char, char) 758 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, ushort, ushort, ushort) 759 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, short, short, short) 760 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, int, int, int) 761 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uint, uint, uint) 762 763 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP 764 765 // binary function (vec & vec) 766 767 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(func_name, func, input_type, output_type) \ 768 __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, const input_type ## 1 & b) \ 769 { \ 770 return VecTraits<output_type ## 1>::make(func (a.x, b.x)); \ 771 } \ 772 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, const input_type ## 2 & b) \ 773 { \ 774 return VecTraits<output_type ## 2>::make(func (a.x, b.x), func (a.y, b.y)); \ 775 } \ 776 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, const input_type ## 3 & b) \ 777 { \ 778 return VecTraits<output_type ## 3>::make(func (a.x, b.x), func (a.y, b.y), func (a.z, b.z)); \ 779 } \ 780 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, const input_type ## 4 & b) \ 781 { \ 782 return VecTraits<output_type ## 4>::make(func (a.x, b.x), func (a.y, b.y), func (a.z, b.z), func (a.w, b.w)); \ 783 } 784 785 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, uchar, uchar) 786 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, char, char) 787 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, ushort, ushort) 788 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, short, short) 789 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, uint, uint) 790 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, int, int) 791 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::fmaxf, float, float) 792 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::fmax, double, double) 793 794 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, uchar, uchar) 795 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, char, char) 796 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, ushort, ushort) 797 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, short, short) 798 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, uint, uint) 799 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, int, int) 800 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::fminf, float, float) 801 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::fmin, double, double) 802 803 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, uchar, float) 804 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, char, float) 805 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, ushort, float) 806 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, short, float) 807 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, uint, float) 808 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, int, float) 809 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, float, float) 810 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypot, double, double) 811 812 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, uchar, float) 813 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, char, float) 814 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, ushort, float) 815 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, short, float) 816 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, uint, float) 817 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, int, float) 818 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, float, float) 819 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2, double, double) 820 821 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC 822 823 // binary function (vec & scalar) 824 825 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(func_name, func, input_type, scalar_type, output_type) \ 826 __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, scalar_type s) \ 827 { \ 828 return VecTraits<output_type ## 1>::make(func ((output_type) a.x, (output_type) s)); \ 829 } \ 830 __device__ __forceinline__ output_type ## 1 func_name(scalar_type s, const input_type ## 1 & b) \ 831 { \ 832 return VecTraits<output_type ## 1>::make(func ((output_type) s, (output_type) b.x)); \ 833 } \ 834 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, scalar_type s) \ 835 { \ 836 return VecTraits<output_type ## 2>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s)); \ 837 } \ 838 __device__ __forceinline__ output_type ## 2 func_name(scalar_type s, const input_type ## 2 & b) \ 839 { \ 840 return VecTraits<output_type ## 2>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y)); \ 841 } \ 842 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, scalar_type s) \ 843 { \ 844 return VecTraits<output_type ## 3>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s), func ((output_type) a.z, (output_type) s)); \ 845 } \ 846 __device__ __forceinline__ output_type ## 3 func_name(scalar_type s, const input_type ## 3 & b) \ 847 { \ 848 return VecTraits<output_type ## 3>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y), func ((output_type) s, (output_type) b.z)); \ 849 } \ 850 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, scalar_type s) \ 851 { \ 852 return VecTraits<output_type ## 4>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s), func ((output_type) a.z, (output_type) s), func ((output_type) a.w, (output_type) s)); \ 853 } \ 854 __device__ __forceinline__ output_type ## 4 func_name(scalar_type s, const input_type ## 4 & b) \ 855 { \ 856 return VecTraits<output_type ## 4>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y), func ((output_type) s, (output_type) b.z), func ((output_type) s, (output_type) b.w)); \ 857 } 858 859 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, uchar, uchar, uchar) 860 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, uchar, float, float) 861 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, uchar, double, double) 862 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, char, char, char) 863 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, char, float, float) 864 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, char, double, double) 865 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, ushort, ushort, ushort) 866 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, ushort, float, float) 867 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, ushort, double, double) 868 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, short, short, short) 869 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, short, float, float) 870 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, short, double, double) 871 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, uint, uint, uint) 872 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, uint, float, float) 873 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, uint, double, double) 874 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, int, int, int) 875 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, int, float, float) 876 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, int, double, double) 877 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, float, float, float) 878 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, float, double, double) 879 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, double, double, double) 880 881 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, uchar, uchar, uchar) 882 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, uchar, float, float) 883 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, uchar, double, double) 884 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, char, char, char) 885 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, char, float, float) 886 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, char, double, double) 887 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, ushort, ushort, ushort) 888 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, ushort, float, float) 889 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, ushort, double, double) 890 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, short, short, short) 891 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, short, float, float) 892 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, short, double, double) 893 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, uint, uint, uint) 894 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, uint, float, float) 895 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, uint, double, double) 896 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, int, int, int) 897 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, int, float, float) 898 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, int, double, double) 899 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, float, float, float) 900 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, float, double, double) 901 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, double, double, double) 902 903 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, uchar, float, float) 904 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, uchar, double, double) 905 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, char, float, float) 906 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, char, double, double) 907 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, ushort, float, float) 908 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, ushort, double, double) 909 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, short, float, float) 910 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, short, double, double) 911 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, uint, float, float) 912 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, uint, double, double) 913 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, int, float, float) 914 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, int, double, double) 915 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, float, float, float) 916 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, float, double, double) 917 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, double, double, double) 918 919 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, uchar, float, float) 920 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, uchar, double, double) 921 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, char, float, float) 922 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, char, double, double) 923 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, ushort, float, float) 924 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, ushort, double, double) 925 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, short, float, float) 926 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, short, double, double) 927 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, uint, float, float) 928 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, uint, double, double) 929 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, int, float, float) 930 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, int, double, double) 931 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, float, float, float) 932 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, float, double, double) 933 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, double, double, double) 934 935 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC 936 937 //! @} 938 939 }} 940 941 #endif 942
