vec_math.hpp
Go to the documentation of this file.
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_GPU_VECMATH_HPP__
44 #define __OPENCV_GPU_VECMATH_HPP__
45 
46 #include "vec_traits.hpp"
47 #include "saturate_cast.hpp"
48 
49 namespace cv { namespace gpu { namespace device
50 {
51 
52 // saturate_cast
53 
54 namespace vec_math_detail
55 {
56  template <int cn, typename VecD> struct SatCastHelper;
57  template <typename VecD> struct SatCastHelper<1, VecD>
58  {
59  template <typename VecS> static __device__ __forceinline__ VecD cast(const VecS& v)
60  {
61  typedef typename VecTraits<VecD>::elem_type D;
62  return VecTraits<VecD>::make(saturate_cast<D>(v.x));
63  }
64  };
65  template <typename VecD> struct SatCastHelper<2, VecD>
66  {
67  template <typename VecS> static __device__ __forceinline__ VecD cast(const VecS& v)
68  {
69  typedef typename VecTraits<VecD>::elem_type D;
70  return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y));
71  }
72  };
73  template <typename VecD> struct SatCastHelper<3, VecD>
74  {
75  template <typename VecS> static __device__ __forceinline__ VecD cast(const VecS& v)
76  {
77  typedef typename VecTraits<VecD>::elem_type D;
78  return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y), saturate_cast<D>(v.z));
79  }
80  };
81  template <typename VecD> struct SatCastHelper<4, VecD>
82  {
83  template <typename VecS> static __device__ __forceinline__ 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), saturate_cast<D>(v.w));
87  }
88  };
89 
90  template <typename VecD, typename VecS> static __device__ __forceinline__ VecD saturate_cast_helper(const VecS& v)
91  {
92  return SatCastHelper<VecTraits<VecD>::cn, VecD>::cast(v);
93  }
94 }
95 
96 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
97 template<typename T> static __device__ __forceinline__ T saturate_cast(const char1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
98 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
99 template<typename T> static __device__ __forceinline__ T saturate_cast(const short1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
100 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
101 template<typename T> static __device__ __forceinline__ T saturate_cast(const int1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
102 template<typename T> static __device__ __forceinline__ T saturate_cast(const float1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
103 template<typename T> static __device__ __forceinline__ T saturate_cast(const double1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
104 
105 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
106 template<typename T> static __device__ __forceinline__ T saturate_cast(const char2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
107 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
108 template<typename T> static __device__ __forceinline__ T saturate_cast(const short2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
109 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
110 template<typename T> static __device__ __forceinline__ T saturate_cast(const int2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
111 template<typename T> static __device__ __forceinline__ T saturate_cast(const float2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
112 template<typename T> static __device__ __forceinline__ T saturate_cast(const double2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
113 
114 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
115 template<typename T> static __device__ __forceinline__ T saturate_cast(const char3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
116 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
117 template<typename T> static __device__ __forceinline__ T saturate_cast(const short3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
118 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
119 template<typename T> static __device__ __forceinline__ T saturate_cast(const int3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
120 template<typename T> static __device__ __forceinline__ T saturate_cast(const float3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
121 template<typename T> static __device__ __forceinline__ T saturate_cast(const double3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
122 
123 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
124 template<typename T> static __device__ __forceinline__ T saturate_cast(const char4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
125 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
126 template<typename T> static __device__ __forceinline__ T saturate_cast(const short4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
127 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
128 template<typename T> static __device__ __forceinline__ T saturate_cast(const int4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
129 template<typename T> static __device__ __forceinline__ T saturate_cast(const float4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
130 template<typename T> static __device__ __forceinline__ T saturate_cast(const double4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
131 
132 // unary operators
133 
134 #define CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(op, input_type, output_type) \
135  __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a) \
136  { \
137  return VecTraits<output_type ## 1>::make(op (a.x)); \
138  } \
139  __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a) \
140  { \
141  return VecTraits<output_type ## 2>::make(op (a.x), op (a.y)); \
142  } \
143  __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a) \
144  { \
145  return VecTraits<output_type ## 3>::make(op (a.x), op (a.y), op (a.z)); \
146  } \
147  __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a) \
148  { \
149  return VecTraits<output_type ## 4>::make(op (a.x), op (a.y), op (a.z), op (a.w)); \
150  }
151 
152 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, char, char)
155 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, float, float)
156 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, double, double)
157 
158 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, uchar, uchar)
159 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, char, uchar)
160 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, ushort, uchar)
161 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, short, uchar)
162 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, int, uchar)
163 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, uint, uchar)
164 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, float, uchar)
165 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, double, uchar)
166 
167 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, uchar, uchar)
168 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, char, char)
169 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, ushort, ushort)
170 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, short, short)
171 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, int, int)
172 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, uint, uint)
173 
174 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_OP
175 
176 // unary functions
177 
178 #define CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(func_name, func, input_type, output_type) \
179  __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a) \
180  { \
181  return VecTraits<output_type ## 1>::make(func (a.x)); \
182  } \
183  __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a) \
184  { \
185  return VecTraits<output_type ## 2>::make(func (a.x), func (a.y)); \
186  } \
187  __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a) \
188  { \
189  return VecTraits<output_type ## 3>::make(func (a.x), func (a.y), func (a.z)); \
190  } \
191  __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a) \
192  { \
193  return VecTraits<output_type ## 4>::make(func (a.x), func (a.y), func (a.z), func (a.w)); \
194  }
195 
196 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, uchar, uchar)
197 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::abs, char, char)
198 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, ushort, ushort)
199 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::abs, short, short)
200 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::abs, int, int)
201 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, uint, uint)
202 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::fabsf, float, float)
203 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::fabs, double, double)
204 
205 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, uchar, float)
206 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, char, float)
207 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, ushort, float)
208 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, short, float)
209 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, int, float)
210 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, uint, float)
211 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, float, float)
212 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrt, double, double)
213 
214 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, uchar, float)
215 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, char, float)
216 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, ushort, float)
217 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, short, float)
218 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, int, float)
219 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, uint, float)
220 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, float, float)
221 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::exp, double, double)
222 
223 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, uchar, float)
224 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, char, float)
225 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, ushort, float)
226 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, short, float)
227 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, int, float)
228 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, uint, float)
229 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, float, float)
230 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2, double, double)
231 
232 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, uchar, float)
233 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, char, float)
234 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, ushort, float)
235 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, short, float)
236 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, int, float)
237 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, uint, float)
238 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, float, float)
239 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10, double, double)
240 
241 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, uchar, float)
242 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, char, float)
243 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, ushort, float)
244 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, short, float)
245 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, int, float)
246 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, uint, float)
247 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, float, float)
248 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::log, double, double)
249 
250 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, uchar, float)
251 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, char, float)
252 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, ushort, float)
253 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, short, float)
254 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, int, float)
255 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, uint, float)
256 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, float, float)
257 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2, double, double)
258 
259 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, uchar, float)
260 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, char, float)
261 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, ushort, float)
262 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, short, float)
263 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, int, float)
264 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, uint, float)
265 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, float, float)
266 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10, double, double)
267 
268 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, uchar, float)
269 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, char, float)
270 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, ushort, float)
271 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, short, float)
272 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, int, float)
273 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, uint, float)
274 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, float, float)
275 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sin, double, double)
276 
277 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, uchar, float)
278 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, char, float)
279 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, ushort, float)
280 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, short, float)
281 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, int, float)
282 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, uint, float)
283 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, float, float)
284 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cos, double, double)
285 
286 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, uchar, float)
287 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, char, float)
288 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, ushort, float)
289 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, short, float)
290 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, int, float)
291 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, uint, float)
292 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, float, float)
293 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tan, double, double)
294 
295 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, uchar, float)
296 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, char, float)
297 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, ushort, float)
298 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, short, float)
299 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, int, float)
300 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, uint, float)
301 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, float, float)
302 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asin, double, double)
303 
304 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, uchar, float)
305 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, char, float)
306 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, ushort, float)
307 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, short, float)
308 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, int, float)
309 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, uint, float)
310 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, float, float)
311 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acos, double, double)
312 
313 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, uchar, float)
314 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, char, float)
315 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, ushort, float)
316 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, short, float)
317 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, int, float)
318 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, uint, float)
319 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, float, float)
320 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atan, double, double)
321 
322 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, uchar, float)
323 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, char, float)
324 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, ushort, float)
325 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, short, float)
326 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, int, float)
327 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, uint, float)
328 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, float, float)
329 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinh, double, double)
330 
331 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, uchar, float)
332 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, char, float)
333 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, ushort, float)
334 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, short, float)
335 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, int, float)
336 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, uint, float)
337 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, float, float)
338 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::cosh, double, double)
339 
340 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, uchar, float)
341 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, char, float)
342 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, ushort, float)
343 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, short, float)
344 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, int, float)
345 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, uint, float)
346 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, float, float)
347 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanh, double, double)
348 
349 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, uchar, float)
350 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, char, float)
351 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, ushort, float)
352 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, short, float)
353 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, int, float)
354 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, uint, float)
355 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, float, float)
356 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinh, double, double)
357 
358 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, uchar, float)
359 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, char, float)
360 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, ushort, float)
361 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, short, float)
362 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, int, float)
363 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, uint, float)
364 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, float, float)
365 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acosh, double, double)
366 
367 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, uchar, float)
368 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, char, float)
369 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, ushort, float)
370 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, short, float)
371 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, int, float)
372 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, uint, float)
373 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, float, float)
374 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanh, double, double)
375 
376 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC
377 
378 // binary operators (vec & vec)
379 
380 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(op, input_type, output_type) \
381  __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, const input_type ## 1 & b) \
382  { \
383  return VecTraits<output_type ## 1>::make(a.x op b.x); \
384  } \
385  __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, const input_type ## 2 & b) \
386  { \
387  return VecTraits<output_type ## 2>::make(a.x op b.x, a.y op b.y); \
388  } \
389  __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, const input_type ## 3 & b) \
390  { \
391  return VecTraits<output_type ## 3>::make(a.x op b.x, a.y op b.y, a.z op b.z); \
392  } \
393  __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, const input_type ## 4 & b) \
394  { \
395  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); \
396  }
397 
400 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, ushort, int)
401 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, short, int)
402 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, int, int)
403 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, uint, uint)
404 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, float, float)
405 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, double, double)
406 
407 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uchar, int)
408 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, char, int)
409 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, ushort, int)
410 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, short, int)
411 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, int, int)
412 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uint, uint)
413 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, float, float)
414 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, double, double)
415 
416 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, uchar, int)
417 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, char, int)
418 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, ushort, int)
419 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, short, int)
420 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, int, int)
421 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, uint, uint)
422 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, float, float)
423 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, double, double)
424 
425 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, uchar, int)
426 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, char, int)
427 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, ushort, int)
428 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, short, int)
429 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, int, int)
430 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, uint, uint)
431 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, float, float)
432 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, double, double)
433 
434 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uchar, uchar)
435 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, char, uchar)
436 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, ushort, uchar)
437 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, short, uchar)
438 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, int, uchar)
439 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uint, uchar)
440 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, float, uchar)
441 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, double, uchar)
442 
443 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uchar, uchar)
444 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, char, uchar)
445 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, ushort, uchar)
446 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, short, uchar)
447 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, int, uchar)
448 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uint, uchar)
449 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, float, uchar)
450 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, double, uchar)
451 
452 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, uchar, uchar)
453 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, char, uchar)
454 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, ushort, uchar)
455 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, short, uchar)
456 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, int, uchar)
457 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, uint, uchar)
458 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, float, uchar)
459 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, double, uchar)
460 
461 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, uchar, uchar)
462 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, char, uchar)
463 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, ushort, uchar)
464 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, short, uchar)
465 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, int, uchar)
466 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, uint, uchar)
467 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, float, uchar)
468 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, double, uchar)
469 
470 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, uchar, uchar)
471 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, char, uchar)
472 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, ushort, uchar)
473 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, short, uchar)
474 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, int, uchar)
475 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, uint, uchar)
476 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, float, uchar)
477 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, double, uchar)
478 
479 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, uchar, uchar)
480 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, char, uchar)
481 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, ushort, uchar)
482 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, short, uchar)
483 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, int, uchar)
484 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, uint, uchar)
485 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, float, uchar)
486 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, double, uchar)
487 
488 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, uchar, uchar)
489 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, char, uchar)
490 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, ushort, uchar)
491 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, short, uchar)
492 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, int, uchar)
493 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, uint, uchar)
494 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, float, uchar)
495 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, double, uchar)
496 
497 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, uchar, uchar)
498 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, char, uchar)
499 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, ushort, uchar)
500 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, short, uchar)
501 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, int, uchar)
502 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, uint, uchar)
503 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, float, uchar)
504 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, double, uchar)
505 
506 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, uchar, uchar)
507 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, char, char)
508 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, ushort, ushort)
509 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, short, short)
510 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, int, int)
511 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, uint, uint)
512 
513 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, uchar, uchar)
514 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, char, char)
515 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, ushort, ushort)
516 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, short, short)
517 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, int, int)
518 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, uint, uint)
519 
520 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uchar, uchar)
521 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, char, char)
522 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, ushort, ushort)
523 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, short, short)
524 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, int, int)
525 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uint, uint)
526 
527 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_OP
528 
529 // binary operators (vec & scalar)
530 
531 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(op, input_type, scalar_type, output_type) \
532  __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, scalar_type s) \
533  { \
534  return VecTraits<output_type ## 1>::make(a.x op s); \
535  } \
536  __device__ __forceinline__ output_type ## 1 operator op(scalar_type s, const input_type ## 1 & b) \
537  { \
538  return VecTraits<output_type ## 1>::make(s op b.x); \
539  } \
540  __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, scalar_type s) \
541  { \
542  return VecTraits<output_type ## 2>::make(a.x op s, a.y op s); \
543  } \
544  __device__ __forceinline__ output_type ## 2 operator op(scalar_type s, const input_type ## 2 & b) \
545  { \
546  return VecTraits<output_type ## 2>::make(s op b.x, s op b.y); \
547  } \
548  __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, scalar_type s) \
549  { \
550  return VecTraits<output_type ## 3>::make(a.x op s, a.y op s, a.z op s); \
551  } \
552  __device__ __forceinline__ output_type ## 3 operator op(scalar_type s, const input_type ## 3 & b) \
553  { \
554  return VecTraits<output_type ## 3>::make(s op b.x, s op b.y, s op b.z); \
555  } \
556  __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, scalar_type s) \
557  { \
558  return VecTraits<output_type ## 4>::make(a.x op s, a.y op s, a.z op s, a.w op s); \
559  } \
560  __device__ __forceinline__ output_type ## 4 operator op(scalar_type s, const input_type ## 4 & b) \
561  { \
562  return VecTraits<output_type ## 4>::make(s op b.x, s op b.y, s op b.z, s op b.w); \
563  }
564 
565 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, int, int)
566 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, float, float)
567 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, double, double)
568 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, int, int)
569 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, float, float)
570 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, double, double)
571 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, int, int)
572 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, float, float)
573 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, double, double)
574 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, int, int)
575 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, float, float)
576 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, double, double)
577 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, int, int)
578 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, float, float)
579 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, double, double)
580 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, uint, uint)
581 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, float, float)
582 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, double, double)
583 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, float, float, float)
584 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, float, double, double)
585 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, double, double, double)
586 
587 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, int, int)
588 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, float, float)
589 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, double, double)
590 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, int, int)
591 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, float, float)
592 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, double, double)
593 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, int, int)
594 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, float, float)
595 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, double, double)
596 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, int, int)
597 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, float, float)
598 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, double, double)
599 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, int, int)
600 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, float, float)
601 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, double, double)
602 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, uint, uint)
603 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, float, float)
604 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, double, double)
605 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, float, float, float)
606 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, float, double, double)
607 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, double, double, double)
608 
609 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, int, int)
610 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, float, float)
611 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, double, double)
612 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, int, int)
613 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, float, float)
614 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, double, double)
615 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, int, int)
616 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, float, float)
617 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, double, double)
618 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, int, int)
619 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, float, float)
620 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, double, double)
621 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, int, int)
622 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, float, float)
623 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, double, double)
624 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, uint, uint)
625 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, float, float)
626 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, double, double)
627 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, float, float, float)
628 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, float, double, double)
629 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, double, double, double)
630 
631 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, int, int)
632 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, float, float)
633 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, double, double)
634 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, int, int)
635 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, float, float)
636 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, double, double)
637 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, int, int)
638 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, float, float)
639 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, double, double)
640 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, int, int)
641 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, float, float)
642 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, double, double)
643 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, int, int)
644 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, float, float)
645 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, double, double)
646 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, uint, uint)
647 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, float, float)
648 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, double, double)
649 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, float, float, float)
650 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, float, double, double)
651 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, double, double, double)
652 
653 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uchar, uchar, uchar)
654 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, char, char, uchar)
655 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, ushort, ushort, uchar)
656 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, short, short, uchar)
657 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, int, int, uchar)
658 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uint, uint, uchar)
659 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, float, float, uchar)
660 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, double, double, uchar)
661 
662 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uchar, uchar, uchar)
663 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, char, char, uchar)
664 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, ushort, ushort, uchar)
665 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, short, short, uchar)
666 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, int, int, uchar)
667 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uint, uint, uchar)
668 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, float, float, uchar)
669 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, double, double, uchar)
670 
671 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, uchar, uchar, uchar)
672 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, char, char, uchar)
673 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, ushort, ushort, uchar)
674 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, short, short, uchar)
675 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, int, int, uchar)
676 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, uint, uint, uchar)
677 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, float, float, uchar)
678 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, double, double, uchar)
679 
680 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, uchar, uchar, uchar)
681 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, char, char, uchar)
682 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, ushort, ushort, uchar)
683 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, short, short, uchar)
684 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, int, int, uchar)
685 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, uint, uint, uchar)
686 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, float, float, uchar)
687 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, double, double, uchar)
688 
689 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, uchar, uchar, uchar)
690 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, char, char, uchar)
691 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, ushort, ushort, uchar)
692 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, short, short, uchar)
693 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, int, int, uchar)
694 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, uint, uint, uchar)
695 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, float, float, uchar)
696 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, double, double, uchar)
697 
698 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, uchar, uchar, uchar)
699 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, char, char, uchar)
700 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, ushort, ushort, uchar)
701 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, short, short, uchar)
702 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, int, int, uchar)
703 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, uint, uint, uchar)
704 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, float, float, uchar)
705 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, double, double, uchar)
706 
707 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, uchar, uchar, uchar)
708 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, char, char, uchar)
709 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, ushort, ushort, uchar)
710 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, short, short, uchar)
711 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, int, int, uchar)
712 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, uint, uint, uchar)
713 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, float, float, uchar)
714 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, double, double, uchar)
715 
716 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, uchar, uchar, uchar)
717 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, char, char, uchar)
718 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, ushort, ushort, uchar)
719 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, short, short, uchar)
720 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, int, int, uchar)
721 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, uint, uint, uchar)
722 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, float, float, uchar)
723 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, double, double, uchar)
724 
725 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, uchar, uchar, uchar)
726 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, char, char, char)
727 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, ushort, ushort, ushort)
728 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, short, short, short)
729 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, int, int, int)
730 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, uint, uint, uint)
731 
732 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, uchar, uchar, uchar)
733 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, char, char, char)
734 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, ushort, ushort, ushort)
735 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, short, short, short)
736 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, int, int, int)
737 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, uint, uint, uint)
738 
739 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uchar, uchar, uchar)
740 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, char, char, char)
741 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, ushort, ushort, ushort)
742 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, short, short, short)
743 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, int, int, int)
744 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uint, uint, uint)
745 
746 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP
747 
748 // binary function (vec & vec)
749 
750 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(func_name, func, input_type, output_type) \
751  __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, const input_type ## 1 & b) \
752  { \
753  return VecTraits<output_type ## 1>::make(func (a.x, b.x)); \
754  } \
755  __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, const input_type ## 2 & b) \
756  { \
757  return VecTraits<output_type ## 2>::make(func (a.x, b.x), func (a.y, b.y)); \
758  } \
759  __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, const input_type ## 3 & b) \
760  { \
761  return VecTraits<output_type ## 3>::make(func (a.x, b.x), func (a.y, b.y), func (a.z, b.z)); \
762  } \
763  __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, const input_type ## 4 & b) \
764  { \
765  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)); \
766  }
767 
769 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, char, char)
770 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, ushort, ushort)
771 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, short, short)
772 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, uint, uint)
773 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, int, int)
774 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::fmaxf, float, float)
775 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::fmax, double, double)
776 
777 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, uchar, uchar)
778 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, char, char)
779 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, ushort, ushort)
780 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, short, short)
781 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, uint, uint)
782 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, int, int)
783 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::fminf, float, float)
784 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::fmin, double, double)
785 
786 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, uchar, float)
787 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, char, float)
788 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, ushort, float)
789 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, short, float)
790 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, uint, float)
791 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, int, float)
792 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, float, float)
793 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypot, double, double)
794 
795 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, uchar, float)
796 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, char, float)
797 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, ushort, float)
798 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, short, float)
799 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, uint, float)
800 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, int, float)
801 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, float, float)
802 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2, double, double)
803 
804 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC
805 
806 // binary function (vec & scalar)
807 
808 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(func_name, func, input_type, scalar_type, output_type) \
809  __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, scalar_type s) \
810  { \
811  return VecTraits<output_type ## 1>::make(func ((output_type) a.x, (output_type) s)); \
812  } \
813  __device__ __forceinline__ output_type ## 1 func_name(scalar_type s, const input_type ## 1 & b) \
814  { \
815  return VecTraits<output_type ## 1>::make(func ((output_type) s, (output_type) b.x)); \
816  } \
817  __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, scalar_type s) \
818  { \
819  return VecTraits<output_type ## 2>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s)); \
820  } \
821  __device__ __forceinline__ output_type ## 2 func_name(scalar_type s, const input_type ## 2 & b) \
822  { \
823  return VecTraits<output_type ## 2>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y)); \
824  } \
825  __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, scalar_type s) \
826  { \
827  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)); \
828  } \
829  __device__ __forceinline__ output_type ## 3 func_name(scalar_type s, const input_type ## 3 & b) \
830  { \
831  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)); \
832  } \
833  __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, scalar_type s) \
834  { \
835  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)); \
836  } \
837  __device__ __forceinline__ output_type ## 4 func_name(scalar_type s, const input_type ## 4 & b) \
838  { \
839  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)); \
840  }
841 
842 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, uchar, uchar, uchar)
843 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, uchar, float, float)
844 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, uchar, double, double)
845 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, char, char, char)
846 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, char, float, float)
847 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, char, double, double)
848 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, ushort, ushort, ushort)
849 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, ushort, float, float)
850 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, ushort, double, double)
851 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, short, short, short)
852 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, short, float, float)
853 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, short, double, double)
854 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, uint, uint, uint)
855 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, uint, float, float)
856 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, uint, double, double)
857 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, int, int, int)
858 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, int, float, float)
859 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, int, double, double)
860 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, float, float, float)
861 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, float, double, double)
862 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, double, double, double)
863 
864 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, uchar, uchar, uchar)
865 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, uchar, float, float)
866 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, uchar, double, double)
867 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, char, char, char)
868 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, char, float, float)
869 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, char, double, double)
870 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, ushort, ushort, ushort)
871 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, ushort, float, float)
872 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, ushort, double, double)
873 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, short, short, short)
874 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, short, float, float)
875 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, short, double, double)
876 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, uint, uint, uint)
877 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, uint, float, float)
878 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, uint, double, double)
879 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, int, int, int)
880 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, int, float, float)
881 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, int, double, double)
882 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, float, float, float)
883 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, float, double, double)
884 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, double, double, double)
885 
886 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, uchar, float, float)
887 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, uchar, double, double)
888 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, char, float, float)
889 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, char, double, double)
890 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, ushort, float, float)
891 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, ushort, double, double)
892 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, short, float, float)
893 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, short, double, double)
894 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, uint, float, float)
895 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, uint, double, double)
896 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, int, float, float)
897 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, int, double, double)
898 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, float, float, float)
899 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, float, double, double)
900 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, double, double, double)
901 
902 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, uchar, float, float)
903 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, uchar, double, double)
904 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, char, float, float)
905 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, char, double, double)
906 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, ushort, float, float)
907 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, ushort, double, double)
908 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, short, float, float)
909 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, short, double, double)
910 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, uint, float, float)
911 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, uint, double, double)
912 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, int, float, float)
913 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, int, double, double)
914 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, float, float, float)
915 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, float, double, double)
916 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, double, double, double)
917 
918 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC
919 
920 }}} // namespace cv { namespace gpu { namespace device
921 
922 #endif // __OPENCV_GPU_VECMATH_HPP__
short float uchar uchar uchar uchar uchar ushort int uchar ushort int float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float int CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, char, int) CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+
CV_EXPORTS void sqrt(const GpuMat &src, GpuMat &dst, Stream &stream=Stream::Null())
CV_EXPORTS void exp(const GpuMat &a, GpuMat &b, Stream &stream=Stream::Null())
static __device__ __forceinline__ VecD cast(const VecS &v)
Definition: vec_math.hpp:83
static __device__ __forceinline__ VecD cast(const VecS &v)
Definition: vec_math.hpp:75
CV_EXPORTS void max(const GpuMat &src1, const GpuMat &src2, GpuMat &dst, Stream &stream=Stream::Null())
computes per-element maximum of two arrays (dst = max(src1, src2))
CV_EXPORTS void min(const GpuMat &src1, const GpuMat &src2, GpuMat &dst, Stream &stream=Stream::Null())
computes per-element minimum of two arrays (dst = min(src1, src2))
CV_EXPORTS void abs(const GpuMat &src, GpuMat &dst, Stream &stream=Stream::Null())
CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max,::max, uchar, uchar, uchar) CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max
const CvMat const CvMat const CvMat CvMat CvMat CvMat CvMat CvSize CvMat CvMat * T
Definition: calib3d.hpp:270
CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, int, int) CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+
short float uchar uchar uchar uchar uchar ushort int uchar CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs,::abs, char, char) CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs
unsigned char uchar
Definition: common.hpp:100
static __device__ __forceinline__ VecD cast(const VecS &v)
Definition: vec_math.hpp:67
CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, char, char) CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-
unsigned int uint
Definition: common.hpp:104
const GLdouble * v
unsigned short ushort
Definition: common.hpp:101
__device__ __forceinline__ _Tp saturate_cast(uchar v)
Definition: saturate_cast.hpp:50
CV_EXPORTS void log(const GpuMat &a, GpuMat &b, Stream &stream=Stream::Null())
static __device__ __forceinline__ VecD cast(const VecS &v)
Definition: vec_math.hpp:59
Definition: vec_traits.hpp:160
CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max,::max, uchar, uchar) CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max