EstervQrCode 1.1.1
Library for qr code manipulation
vec_math.hpp
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 // Third party copyrights are property of their respective owners.
16 //
17 // Redistribution and use in source and binary forms, with or without modification,
18 // are permitted provided that the following conditions are met:
19 //
20 // * Redistribution's of source code must retain the above copyright notice,
21 // this list of conditions and the following disclaimer.
22 //
23 // * Redistribution's in binary form must reproduce the above copyright notice,
24 // this list of conditions and the following disclaimer in the documentation
25 // and/or other materials provided with the distribution.
26 //
27 // * The name of the copyright holders may not be used to endorse or promote products
28 // derived from this software without specific prior written permission.
29 //
30 // This software is provided by the copyright holders and contributors "as is" and
31 // any express or implied warranties, including, but not limited to, the implied
32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
33 // In no event shall the Intel Corporation or contributors be liable for any direct,
34 // indirect, incidental, special, exemplary, or consequential damages
35 // (including, but not limited to, procurement of substitute goods or services;
36 // loss of use, data, or profits; or business interruption) however caused
37 // and on any theory of liability, whether in contract, strict liability,
38 // or tort (including negligence or otherwise) arising in any way out of
39 // the use of this software, even if advised of the possibility of such damage.
40 //
41 //M*/
42 
43 #ifndef OPENCV_CUDA_VECMATH_HPP
44 #define OPENCV_CUDA_VECMATH_HPP
45 
46 #include "vec_traits.hpp"
47 #include "saturate_cast.hpp"
48 
54 
55 namespace cv { namespace cuda { namespace device
56 {
57 
58 // saturate_cast
59 
60 namespace vec_math_detail
61 {
62  template <int cn, typename VecD> struct SatCastHelper;
63  template <typename VecD> struct SatCastHelper<1, VecD>
64  {
65  template <typename VecS> static __device__ __forceinline__ 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  template <typename VecD> struct SatCastHelper<2, VecD>
72  {
73  template <typename VecS> static __device__ __forceinline__ VecD cast(const VecS& v)
74  {
75  typedef typename VecTraits<VecD>::elem_type D;
76  return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y));
77  }
78  };
79  template <typename VecD> struct SatCastHelper<3, VecD>
80  {
81  template <typename VecS> static __device__ __forceinline__ VecD cast(const VecS& v)
82  {
83  typedef typename VecTraits<VecD>::elem_type D;
84  return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y), saturate_cast<D>(v.z));
85  }
86  };
87  template <typename VecD> struct SatCastHelper<4, VecD>
88  {
89  template <typename VecS> static __device__ __forceinline__ VecD cast(const VecS& v)
90  {
91  typedef typename VecTraits<VecD>::elem_type D;
92  return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y), saturate_cast<D>(v.z), saturate_cast<D>(v.w));
93  }
94  };
95 
96  template <typename VecD, typename VecS> static __device__ __forceinline__ VecD saturate_cast_helper(const VecS& v)
97  {
98  return SatCastHelper<VecTraits<VecD>::cn, VecD>::cast(v);
99  }
100 }
101 
102 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
103 template<typename T> static __device__ __forceinline__ T saturate_cast(const char1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
104 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
105 template<typename T> static __device__ __forceinline__ T saturate_cast(const short1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
106 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
107 template<typename T> static __device__ __forceinline__ T saturate_cast(const int1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
108 template<typename T> static __device__ __forceinline__ T saturate_cast(const float1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
109 template<typename T> static __device__ __forceinline__ T saturate_cast(const double1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
110 
111 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
112 template<typename T> static __device__ __forceinline__ T saturate_cast(const char2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
113 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
114 template<typename T> static __device__ __forceinline__ T saturate_cast(const short2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
115 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
116 template<typename T> static __device__ __forceinline__ T saturate_cast(const int2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
117 template<typename T> static __device__ __forceinline__ T saturate_cast(const float2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
118 template<typename T> static __device__ __forceinline__ T saturate_cast(const double2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
119 
120 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
121 template<typename T> static __device__ __forceinline__ T saturate_cast(const char3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
122 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
123 template<typename T> static __device__ __forceinline__ T saturate_cast(const short3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
124 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
125 template<typename T> static __device__ __forceinline__ T saturate_cast(const int3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
126 template<typename T> static __device__ __forceinline__ T saturate_cast(const float3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
127 template<typename T> static __device__ __forceinline__ T saturate_cast(const double3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
128 
129 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
130 template<typename T> static __device__ __forceinline__ T saturate_cast(const char4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
131 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
132 template<typename T> static __device__ __forceinline__ T saturate_cast(const short4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
133 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
134 template<typename T> static __device__ __forceinline__ T saturate_cast(const int4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
135 template<typename T> static __device__ __forceinline__ T saturate_cast(const float4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
136 template<typename T> static __device__ __forceinline__ T saturate_cast(const double4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
137 
138 // unary operators
139 
140 #define CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(op, input_type, output_type) \
141  __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a) \
142  { \
143  return VecTraits<output_type ## 1>::make(op (a.x)); \
144  } \
145  __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a) \
146  { \
147  return VecTraits<output_type ## 2>::make(op (a.x), op (a.y)); \
148  } \
149  __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a) \
150  { \
151  return VecTraits<output_type ## 3>::make(op (a.x), op (a.y), op (a.z)); \
152  } \
153  __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a) \
154  { \
155  return VecTraits<output_type ## 4>::make(op (a.x), op (a.y), op (a.z), op (a.w)); \
156  }
157 
158 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, char, char)
159 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, short, short)
160 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, int, int)
161 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, float, float)
162 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, double, double)
163 
164 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, uchar, uchar)
165 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, char, uchar)
166 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, ushort, uchar)
167 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, short, uchar)
168 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, int, uchar)
169 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, uint, uchar)
170 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, float, uchar)
171 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, double, uchar)
172 
173 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, uchar, uchar)
174 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, char, char)
175 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, ushort, ushort)
176 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, short, short)
177 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, int, int)
178 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, uint, uint)
179 
180 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_OP
181 
182 // unary functions
183 
184 #define CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(func_name, func, input_type, output_type) \
185  __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a) \
186  { \
187  return VecTraits<output_type ## 1>::make(func (a.x)); \
188  } \
189  __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a) \
190  { \
191  return VecTraits<output_type ## 2>::make(func (a.x), func (a.y)); \
192  } \
193  __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a) \
194  { \
195  return VecTraits<output_type ## 3>::make(func (a.x), func (a.y), func (a.z)); \
196  } \
197  __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a) \
198  { \
199  return VecTraits<output_type ## 4>::make(func (a.x), func (a.y), func (a.z), func (a.w)); \
200  }
201 
202 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::fabsf, float, float)
203 
204 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, uchar, float)
205 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, char, float)
206 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, ushort, float)
207 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, short, float)
208 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, int, float)
209 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, uint, float)
210 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, float, float)
211 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrt, double, double)
212 
213 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, uchar, float)
214 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, char, float)
215 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, ushort, float)
216 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, short, float)
217 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, int, float)
218 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, uint, float)
219 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, float, float)
220 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::exp, double, double)
221 
222 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, uchar, float)
223 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, char, float)
224 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, ushort, float)
225 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, short, float)
226 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, int, float)
227 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, uint, float)
228 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, float, float)
229 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2, double, double)
230 
231 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, uchar, float)
232 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, char, float)
233 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, ushort, float)
234 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, short, float)
235 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, int, float)
236 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, uint, float)
237 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, float, float)
238 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10, double, double)
239 
240 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, uchar, float)
241 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, char, float)
242 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, ushort, float)
243 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, short, float)
244 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, int, float)
245 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, uint, float)
246 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, float, float)
247 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::log, double, double)
248 
249 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, uchar, float)
250 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, char, float)
251 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, ushort, float)
252 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, short, float)
253 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, int, float)
254 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, uint, float)
255 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, float, float)
256 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2, double, double)
257 
258 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, uchar, float)
259 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, char, float)
260 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, ushort, float)
261 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, short, float)
262 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, int, float)
263 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, uint, float)
264 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, float, float)
265 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10, double, double)
266 
267 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, uchar, float)
268 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, char, float)
269 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, ushort, float)
270 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, short, float)
271 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, int, float)
272 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, uint, float)
273 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, float, float)
274 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sin, double, double)
275 
276 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, uchar, float)
277 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, char, float)
278 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, ushort, float)
279 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, short, float)
280 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, int, float)
281 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, uint, float)
282 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, float, float)
283 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cos, double, double)
284 
285 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, uchar, float)
286 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, char, float)
287 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, ushort, float)
288 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, short, float)
289 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, int, float)
290 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, uint, float)
291 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, float, float)
292 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tan, double, double)
293 
294 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, uchar, float)
295 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, char, float)
296 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, ushort, float)
297 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, short, float)
298 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, int, float)
299 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, uint, float)
300 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, float, float)
301 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asin, double, double)
302 
303 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, uchar, float)
304 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, char, float)
305 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, ushort, float)
306 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, short, float)
307 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, int, float)
308 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, uint, float)
309 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, float, float)
310 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acos, double, double)
311 
312 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, uchar, float)
313 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, char, float)
314 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, ushort, float)
315 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, short, float)
316 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, int, float)
317 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, uint, float)
318 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, float, float)
319 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atan, double, double)
320 
321 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, uchar, float)
322 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, char, float)
323 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, ushort, float)
324 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, short, float)
325 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, int, float)
326 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, uint, float)
327 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, float, float)
328 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinh, double, double)
329 
330 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, uchar, float)
331 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, char, float)
332 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, ushort, float)
333 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, short, float)
334 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, int, float)
335 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, uint, float)
336 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, float, float)
337 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::cosh, double, double)
338 
339 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, uchar, float)
340 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, char, float)
341 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, ushort, float)
342 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, short, float)
343 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, int, float)
344 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, uint, float)
345 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, float, float)
346 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanh, double, double)
347 
348 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, uchar, float)
349 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, char, float)
350 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, ushort, float)
351 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, short, float)
352 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, int, float)
353 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, uint, float)
354 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, float, float)
355 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinh, double, double)
356 
357 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, uchar, float)
358 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, char, float)
359 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, ushort, float)
360 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, short, float)
361 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, int, float)
362 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, uint, float)
363 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, float, float)
364 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acosh, double, double)
365 
366 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, uchar, float)
367 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, char, float)
368 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, ushort, float)
369 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, short, float)
370 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, int, float)
371 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, uint, float)
372 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, float, float)
373 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanh, double, double)
374 
375 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC
376 
377 // binary operators (vec & vec)
378 
379 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(op, input_type, output_type) \
380  __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, const input_type ## 1 & b) \
381  { \
382  return VecTraits<output_type ## 1>::make(a.x op b.x); \
383  } \
384  __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, const input_type ## 2 & b) \
385  { \
386  return VecTraits<output_type ## 2>::make(a.x op b.x, a.y op b.y); \
387  } \
388  __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, const input_type ## 3 & b) \
389  { \
390  return VecTraits<output_type ## 3>::make(a.x op b.x, a.y op b.y, a.z op b.z); \
391  } \
392  __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, const input_type ## 4 & b) \
393  { \
394  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); \
395  }
396 
397 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, uchar, int)
398 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, char, int)
399 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, ushort, int)
400 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, short, int)
401 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, int, int)
402 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, uint, uint)
403 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, float, float)
404 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, double, double)
405 
406 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uchar, int)
407 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, char, int)
408 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, ushort, int)
409 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, short, int)
410 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, int, int)
411 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uint, uint)
412 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, float, float)
413 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, double, double)
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, uchar)
434 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, char, uchar)
435 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, ushort, uchar)
436 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, short, uchar)
437 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, int, uchar)
438 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uint, uchar)
439 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, float, uchar)
440 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, double, uchar)
441 
442 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uchar, uchar)
443 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, char, uchar)
444 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, ushort, uchar)
445 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, short, uchar)
446 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, int, uchar)
447 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uint, uchar)
448 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, float, uchar)
449 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, double, uchar)
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, char)
507 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, ushort, ushort)
508 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, short, short)
509 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, int, int)
510 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, uint, uint)
511 
512 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, uchar, uchar)
513 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, char, char)
514 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, ushort, ushort)
515 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, short, short)
516 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, int, int)
517 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, uint, uint)
518 
519 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uchar, uchar)
520 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, char, char)
521 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, ushort, ushort)
522 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, short, short)
523 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, int, int)
524 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uint, uint)
525 
526 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_OP
527 
528 // binary operators (vec & scalar)
529 
530 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(op, input_type, scalar_type, output_type) \
531  __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, scalar_type s) \
532  { \
533  return VecTraits<output_type ## 1>::make(a.x op s); \
534  } \
535  __device__ __forceinline__ output_type ## 1 operator op(scalar_type s, const input_type ## 1 & b) \
536  { \
537  return VecTraits<output_type ## 1>::make(s op b.x); \
538  } \
539  __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, scalar_type s) \
540  { \
541  return VecTraits<output_type ## 2>::make(a.x op s, a.y op s); \
542  } \
543  __device__ __forceinline__ output_type ## 2 operator op(scalar_type s, const input_type ## 2 & b) \
544  { \
545  return VecTraits<output_type ## 2>::make(s op b.x, s op b.y); \
546  } \
547  __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, scalar_type s) \
548  { \
549  return VecTraits<output_type ## 3>::make(a.x op s, a.y op s, a.z op s); \
550  } \
551  __device__ __forceinline__ output_type ## 3 operator op(scalar_type s, const input_type ## 3 & b) \
552  { \
553  return VecTraits<output_type ## 3>::make(s op b.x, s op b.y, s op b.z); \
554  } \
555  __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, scalar_type s) \
556  { \
557  return VecTraits<output_type ## 4>::make(a.x op s, a.y op s, a.z op s, a.w op s); \
558  } \
559  __device__ __forceinline__ output_type ## 4 operator op(scalar_type s, const input_type ## 4 & b) \
560  { \
561  return VecTraits<output_type ## 4>::make(s op b.x, s op b.y, s op b.z, s op b.w); \
562  }
563 
564 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, int, int)
565 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, float, float)
566 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, double, double)
567 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, int, int)
568 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, float, float)
569 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, double, double)
570 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, int, int)
571 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, float, float)
572 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, double, double)
573 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, int, int)
574 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, float, float)
575 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, double, double)
576 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, int, int)
577 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, float, float)
578 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, double, double)
579 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, uint, uint)
580 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, float, float)
581 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, double, double)
582 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, float, float, float)
583 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, float, double, double)
584 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, double, double, double)
585 
586 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, int, int)
587 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, float, float)
588 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, double, double)
589 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, int, int)
590 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, float, float)
591 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, double, double)
592 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, int, int)
593 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, float, float)
594 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, double, double)
595 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, int, int)
596 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, float, float)
597 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, double, double)
598 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, int, int)
599 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, float, float)
600 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, double, double)
601 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, uint, uint)
602 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, float, float)
603 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, double, double)
604 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, float, float, float)
605 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, float, double, double)
606 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, double, double, double)
607 
608 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, int, int)
609 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, float, float)
610 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, double, double)
611 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, int, int)
612 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, float, float)
613 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, double, double)
614 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, int, int)
615 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, float, float)
616 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, double, double)
617 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, int, int)
618 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, float, float)
619 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, double, double)
620 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, int, int)
621 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, float, float)
622 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, double, double)
623 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, uint, uint)
624 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, float, float)
625 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, double, double)
626 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, float, float, float)
627 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, float, double, double)
628 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, double, double, double)
629 
630 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, int, int)
631 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, float, float)
632 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, double, double)
633 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, int, int)
634 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, float, float)
635 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, double, double)
636 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, int, int)
637 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, float, float)
638 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, double, double)
639 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, int, int)
640 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, float, float)
641 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, double, double)
642 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, int, int)
643 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, float, float)
644 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, double, double)
645 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, uint, uint)
646 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, float, float)
647 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, double, double)
648 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, float, float, float)
649 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, float, double, double)
650 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, double, double, double)
651 
652 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uchar, uchar, uchar)
653 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, char, char, uchar)
654 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, ushort, ushort, uchar)
655 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, short, short, uchar)
656 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, int, int, uchar)
657 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uint, uint, uchar)
658 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, float, float, uchar)
659 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, double, double, uchar)
660 
661 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uchar, uchar, uchar)
662 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, char, char, uchar)
663 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, ushort, ushort, uchar)
664 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, short, short, uchar)
665 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, int, int, uchar)
666 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uint, uint, uchar)
667 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, float, float, uchar)
668 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, double, double, uchar)
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, char)
726 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, ushort, ushort, ushort)
727 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, short, short, short)
728 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, int, int, int)
729 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, uint, uint, uint)
730 
731 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, uchar, uchar, uchar)
732 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, char, char, char)
733 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, ushort, ushort, ushort)
734 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, short, short, short)
735 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, int, int, int)
736 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, uint, uint, uint)
737 
738 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uchar, uchar, uchar)
739 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, char, char, char)
740 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, ushort, ushort, ushort)
741 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, short, short, short)
742 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, int, int, int)
743 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uint, uint, uint)
744 
745 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP
746 
747 // binary function (vec & vec)
748 
749 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(func_name, func, input_type, output_type) \
750  __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, const input_type ## 1 & b) \
751  { \
752  return VecTraits<output_type ## 1>::make(func (a.x, b.x)); \
753  } \
754  __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, const input_type ## 2 & b) \
755  { \
756  return VecTraits<output_type ## 2>::make(func (a.x, b.x), func (a.y, b.y)); \
757  } \
758  __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, const input_type ## 3 & b) \
759  { \
760  return VecTraits<output_type ## 3>::make(func (a.x, b.x), func (a.y, b.y), func (a.z, b.z)); \
761  } \
762  __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, const input_type ## 4 & b) \
763  { \
764  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)); \
765  }
766 
767 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, uchar, uchar)
768 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, char, char)
769 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, ushort, ushort)
770 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, short, short)
771 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, uint, uint)
772 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, int, int)
773 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::fmaxf, float, float)
774 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::fmax, double, double)
775 
776 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, uchar, uchar)
777 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, char, char)
778 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, ushort, ushort)
779 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, short, short)
780 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, uint, uint)
781 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, int, int)
782 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::fminf, float, float)
783 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::fmin, double, double)
784 
785 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, uchar, float)
786 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, char, float)
787 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, ushort, float)
788 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, short, float)
789 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, uint, float)
790 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, int, float)
791 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, float, float)
792 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypot, double, double)
793 
794 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, uchar, float)
795 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, char, float)
796 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, ushort, float)
797 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, short, float)
798 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, uint, float)
799 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, int, float)
800 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, float, float)
801 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2, double, double)
802 
803 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC
804 
805 // binary function (vec & scalar)
806 
807 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(func_name, func, input_type, scalar_type, output_type) \
808  __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, scalar_type s) \
809  { \
810  return VecTraits<output_type ## 1>::make(func ((output_type) a.x, (output_type) s)); \
811  } \
812  __device__ __forceinline__ output_type ## 1 func_name(scalar_type s, const input_type ## 1 & b) \
813  { \
814  return VecTraits<output_type ## 1>::make(func ((output_type) s, (output_type) b.x)); \
815  } \
816  __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, scalar_type s) \
817  { \
818  return VecTraits<output_type ## 2>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s)); \
819  } \
820  __device__ __forceinline__ output_type ## 2 func_name(scalar_type s, const input_type ## 2 & b) \
821  { \
822  return VecTraits<output_type ## 2>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y)); \
823  } \
824  __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, scalar_type s) \
825  { \
826  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)); \
827  } \
828  __device__ __forceinline__ output_type ## 3 func_name(scalar_type s, const input_type ## 3 & b) \
829  { \
830  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)); \
831  } \
832  __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, scalar_type s) \
833  { \
834  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)); \
835  } \
836  __device__ __forceinline__ output_type ## 4 func_name(scalar_type s, const input_type ## 4 & b) \
837  { \
838  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)); \
839  }
840 
841 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, uchar, uchar, uchar)
842 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, uchar, float, float)
843 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, uchar, double, double)
844 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, char, char, char)
845 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, char, float, float)
846 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, char, double, double)
847 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, ushort, ushort, ushort)
848 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, ushort, float, float)
849 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, ushort, double, double)
850 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, short, short, short)
851 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, short, float, float)
852 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, short, double, double)
853 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, uint, uint, uint)
854 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, uint, float, float)
855 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, uint, double, double)
856 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, int, int, int)
857 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, int, float, float)
858 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, int, double, double)
859 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, float, float, float)
860 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, float, double, double)
861 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, double, double, double)
862 
863 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, uchar, uchar, uchar)
864 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, uchar, float, float)
865 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, uchar, double, double)
866 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, char, char, char)
867 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, char, float, float)
868 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, char, double, double)
869 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, ushort, ushort, ushort)
870 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, ushort, float, float)
871 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, ushort, double, double)
872 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, short, short, short)
873 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, short, float, float)
874 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, short, double, double)
875 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, uint, uint, uint)
876 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, uint, float, float)
877 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, uint, double, double)
878 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, int, int, int)
879 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, int, float, float)
880 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, int, double, double)
881 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, float, float, float)
882 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, float, double, double)
883 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, double, double, double)
884 
885 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, uchar, float, float)
886 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, uchar, double, double)
887 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, char, float, float)
888 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, char, double, double)
889 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, ushort, float, float)
890 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, ushort, double, double)
891 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, short, float, float)
892 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, short, double, double)
893 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, uint, float, float)
894 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, uint, double, double)
895 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, int, float, float)
896 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, int, double, double)
897 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, float, float, float)
898 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, float, double, double)
899 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, double, double, double)
900 
901 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, uchar, float, float)
902 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, uchar, double, double)
903 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, char, float, float)
904 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, char, double, double)
905 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, ushort, float, float)
906 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, ushort, double, double)
907 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, short, float, float)
908 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, short, double, double)
909 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, uint, float, float)
910 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, uint, double, double)
911 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, int, float, float)
912 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, int, double, double)
913 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, float, float, float)
914 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, float, double, double)
915 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, double, double, double)
916 
917 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC
918 
919 }}} // namespace cv { namespace cuda { namespace device
920 
922 
923 #endif // OPENCV_CUDA_VECMATH_HPP
cv::T
InputArrayOfArrays InputArrayOfArrays InputOutputArray InputOutputArray InputOutputArray InputOutputArray Size InputOutputArray InputOutputArray T
Definition: calib3d.hpp:1867
uint
uint32_t uint
Definition: interface.h:42
uchar
unsigned char uchar
Definition: interface.h:51
ushort
unsigned short ushort
Definition: interface.h:52
cv::max
softfloat max(const softfloat &a, const softfloat &b)
Definition: softfloat.hpp:440
cv::min
softfloat min(const softfloat &a, const softfloat &b)
Min and Max functions.
Definition: softfloat.hpp:437
cv::saturate_cast
static _Tp saturate_cast(uchar v)
Template function for accurate conversion from one primitive type to another.
Definition: saturate.hpp:81
cv::asinh
Quat< T > asinh(const Quat< T > &q)
cv::tanh
Quat< T > tanh(const Quat< T > &q)
cv::acosh
Quat< T > acosh(const Quat< T > &q)
cv::acos
Quat< T > acos(const Quat< T > &q)
cv::tan
Quat< T > tan(const Quat< T > &q)
cv::sin
Quat< T > sin(const Quat< T > &q)
cv::atan
Quat< T > atan(const Quat< T > &q)
cv::asin
Quat< T > asin(const Quat< T > &q)
cv::sinh
Quat< T > sinh(const Quat< T > &q)
cv::sqrt
Quat< S > sqrt(const Quat< S > &q, QuatAssumeType assumeUnit=QUAT_ASSUME_NOT_UNIT)
cv::cosh
Quat< T > cosh(const Quat< T > &q)
cv::atanh
Quat< T > atanh(const Quat< T > &q)
cv::cos
Quat< T > cos(const Quat< T > &q)
cv
"black box" representation of the file storage associated with a file on disk.
Definition: calib3d.hpp:441
cv::log
DualQuat< T > log(const DualQuat< T > &dq, QuatAssumeType assumeUnit=QUAT_ASSUME_NOT_UNIT)
Definition: dualquaternion.inl.hpp:344
cv::exp
DualQuat< T > exp(const DualQuat< T > &dq)
Definition: dualquaternion.inl.hpp:312
cv::abs
static uchar abs(uchar a)
Definition: cvstd.hpp:66