• Home
  • Line#
  • Scopes#
  • Navigate#
  • Raw
  • Download
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 
49 /** @file
50  * @deprecated Use @ref cudev instead.
51  */
52 
53 //! @cond IGNORED
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     {
castcv::cuda::device::vec_math_detail::SatCastHelper65         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     {
castcv::cuda::device::vec_math_detail::SatCastHelper73         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     {
castcv::cuda::device::vec_math_detail::SatCastHelper81         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     {
castcv::cuda::device::vec_math_detail::SatCastHelper89         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 
saturate_cast_helper(const VecS & v)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 
saturate_cast(const uchar1 & v)102 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const char1 & v)103 template<typename T> static __device__ __forceinline__ T saturate_cast(const char1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const ushort1 & v)104 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const short1 & v)105 template<typename T> static __device__ __forceinline__ T saturate_cast(const short1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const uint1 & v)106 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const int1 & v)107 template<typename T> static __device__ __forceinline__ T saturate_cast(const int1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const float1 & v)108 template<typename T> static __device__ __forceinline__ T saturate_cast(const float1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const double1 & v)109 template<typename T> static __device__ __forceinline__ T saturate_cast(const double1& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
110 
saturate_cast(const uchar2 & v)111 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const char2 & v)112 template<typename T> static __device__ __forceinline__ T saturate_cast(const char2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const ushort2 & v)113 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const short2 & v)114 template<typename T> static __device__ __forceinline__ T saturate_cast(const short2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const uint2 & v)115 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const int2 & v)116 template<typename T> static __device__ __forceinline__ T saturate_cast(const int2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const float2 & v)117 template<typename T> static __device__ __forceinline__ T saturate_cast(const float2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const double2 & v)118 template<typename T> static __device__ __forceinline__ T saturate_cast(const double2& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
119 
saturate_cast(const uchar3 & v)120 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const char3 & v)121 template<typename T> static __device__ __forceinline__ T saturate_cast(const char3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const ushort3 & v)122 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const short3 & v)123 template<typename T> static __device__ __forceinline__ T saturate_cast(const short3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const uint3 & v)124 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const int3 & v)125 template<typename T> static __device__ __forceinline__ T saturate_cast(const int3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const float3 & v)126 template<typename T> static __device__ __forceinline__ T saturate_cast(const float3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const double3 & v)127 template<typename T> static __device__ __forceinline__ T saturate_cast(const double3& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
128 
saturate_cast(const uchar4 & v)129 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const char4 & v)130 template<typename T> static __device__ __forceinline__ T saturate_cast(const char4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const ushort4 & v)131 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const short4 & v)132 template<typename T> static __device__ __forceinline__ T saturate_cast(const short4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const uint4 & v)133 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const int4 & v)134 template<typename T> static __device__ __forceinline__ T saturate_cast(const int4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const float4 & v)135 template<typename T> static __device__ __forceinline__ T saturate_cast(const float4& v) {return vec_math_detail::saturate_cast_helper<T>(v);}
saturate_cast(const double4 & 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, /*::abs*/, uchar, uchar)
203 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::abs, char, char)
204 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, ushort, ushort)
205 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::abs, short, short)
206 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::abs, int, int)
207 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, uint, uint)
208 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::fabsf, float, float)
209 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::fabs, double, double)
210 
211 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, uchar, float)
212 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, char, float)
213 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, ushort, float)
214 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, short, float)
215 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, int, float)
216 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, uint, float)
217 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, float, float)
218 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrt, double, double)
219 
220 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, uchar, float)
221 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, char, float)
222 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, ushort, float)
223 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, short, float)
224 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, int, float)
225 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, uint, float)
226 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, float, float)
227 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::exp, double, double)
228 
229 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, uchar, float)
230 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, char, float)
231 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, ushort, float)
232 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, short, float)
233 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, int, float)
234 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, uint, float)
235 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, float, float)
236 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2, double, double)
237 
238 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, uchar, float)
239 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, char, float)
240 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, ushort, float)
241 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, short, float)
242 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, int, float)
243 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, uint, float)
244 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, float, float)
245 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10, double, double)
246 
247 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, uchar, float)
248 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, char, float)
249 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, ushort, float)
250 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, short, float)
251 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, int, float)
252 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, uint, float)
253 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, float, float)
254 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::log, double, double)
255 
256 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, uchar, float)
257 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, char, float)
258 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, ushort, float)
259 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, short, float)
260 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, int, float)
261 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, uint, float)
262 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, float, float)
263 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2, double, double)
264 
265 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, uchar, float)
266 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, char, float)
267 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, ushort, float)
268 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, short, float)
269 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, int, float)
270 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, uint, float)
271 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, float, float)
272 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10, double, double)
273 
274 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, uchar, float)
275 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, char, float)
276 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, ushort, float)
277 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, short, float)
278 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, int, float)
279 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, uint, float)
280 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, float, float)
281 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sin, double, double)
282 
283 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, uchar, float)
284 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, char, float)
285 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, ushort, float)
286 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, short, float)
287 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, int, float)
288 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, uint, float)
289 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, float, float)
290 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cos, double, double)
291 
292 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, uchar, float)
293 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, char, float)
294 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, ushort, float)
295 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, short, float)
296 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, int, float)
297 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, uint, float)
298 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, float, float)
299 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tan, double, double)
300 
301 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, uchar, float)
302 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, char, float)
303 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, ushort, float)
304 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, short, float)
305 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, int, float)
306 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, uint, float)
307 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, float, float)
308 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asin, double, double)
309 
310 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, uchar, float)
311 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, char, float)
312 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, ushort, float)
313 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, short, float)
314 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, int, float)
315 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, uint, float)
316 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, float, float)
317 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acos, double, double)
318 
319 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, uchar, float)
320 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, char, float)
321 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, ushort, float)
322 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, short, float)
323 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, int, float)
324 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, uint, float)
325 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, float, float)
326 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atan, double, double)
327 
328 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, uchar, float)
329 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, char, float)
330 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, ushort, float)
331 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, short, float)
332 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, int, float)
333 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, uint, float)
334 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, float, float)
335 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinh, double, double)
336 
337 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, uchar, float)
338 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, char, float)
339 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, ushort, float)
340 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, short, float)
341 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, int, float)
342 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, uint, float)
343 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, float, float)
344 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::cosh, double, double)
345 
346 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, uchar, float)
347 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, char, float)
348 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, ushort, float)
349 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, short, float)
350 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, int, float)
351 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, uint, float)
352 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, float, float)
353 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanh, double, double)
354 
355 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, uchar, float)
356 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, char, float)
357 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, ushort, float)
358 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, short, float)
359 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, int, float)
360 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, uint, float)
361 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, float, float)
362 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinh, double, double)
363 
364 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, uchar, float)
365 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, char, float)
366 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, ushort, float)
367 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, short, float)
368 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, int, float)
369 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, uint, float)
370 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, float, float)
371 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acosh, double, double)
372 
373 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, uchar, float)
374 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, char, float)
375 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, ushort, float)
376 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, short, float)
377 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, int, float)
378 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, uint, float)
379 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, float, float)
380 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanh, double, double)
381 
382 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC
383 
384 // binary operators (vec & vec)
385 
386 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(op, input_type, output_type) \
387     __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, const input_type ## 1 & b) \
388     { \
389         return VecTraits<output_type ## 1>::make(a.x op b.x); \
390     } \
391     __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, const input_type ## 2 & b) \
392     { \
393         return VecTraits<output_type ## 2>::make(a.x op b.x, a.y op b.y); \
394     } \
395     __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, const input_type ## 3 & b) \
396     { \
397         return VecTraits<output_type ## 3>::make(a.x op b.x, a.y op b.y, a.z op b.z); \
398     } \
399     __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, const input_type ## 4 & b) \
400     { \
401         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); \
402     }
403 
404 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, uchar, int)
405 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, char, int)
406 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, ushort, int)
407 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, short, int)
408 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, int, int)
409 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, uint, uint)
410 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, float, float)
411 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, double, double)
412 
413 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uchar, int)
414 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, char, int)
415 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, ushort, int)
416 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, short, int)
417 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, int, int)
418 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uint, uint)
419 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, float, float)
420 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, double, double)
421 
422 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, uchar, int)
423 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, char, int)
424 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, ushort, int)
425 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, short, int)
426 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, int, int)
427 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, uint, uint)
428 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, float, float)
429 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, double, double)
430 
431 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, uchar, int)
432 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, char, int)
433 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, ushort, int)
434 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, short, int)
435 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, int, int)
436 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, uint, uint)
437 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, float, float)
438 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, double, double)
439 
440 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uchar, uchar)
441 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, char, uchar)
442 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, ushort, uchar)
443 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, short, uchar)
444 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, int, uchar)
445 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uint, uchar)
446 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, float, uchar)
447 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, double, uchar)
448 
449 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uchar, uchar)
450 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, char, uchar)
451 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, ushort, uchar)
452 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, short, uchar)
453 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, int, uchar)
454 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uint, uchar)
455 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, float, uchar)
456 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, double, uchar)
457 
458 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, uchar, uchar)
459 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, char, uchar)
460 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, ushort, uchar)
461 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, short, uchar)
462 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, int, uchar)
463 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, uint, uchar)
464 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, float, uchar)
465 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, double, uchar)
466 
467 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, uchar, uchar)
468 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, char, uchar)
469 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, ushort, uchar)
470 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, short, uchar)
471 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, int, uchar)
472 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, uint, uchar)
473 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, float, uchar)
474 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, double, uchar)
475 
476 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, uchar, uchar)
477 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, char, uchar)
478 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, ushort, uchar)
479 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, short, uchar)
480 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, int, uchar)
481 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, uint, uchar)
482 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, float, uchar)
483 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, double, uchar)
484 
485 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, uchar, uchar)
486 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, char, uchar)
487 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, ushort, uchar)
488 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, short, uchar)
489 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, int, uchar)
490 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, uint, uchar)
491 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, float, uchar)
492 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, double, uchar)
493 
494 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, uchar, uchar)
495 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, char, uchar)
496 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, ushort, uchar)
497 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, short, uchar)
498 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, int, uchar)
499 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, uint, uchar)
500 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, float, uchar)
501 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, double, uchar)
502 
503 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, uchar, uchar)
504 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, char, uchar)
505 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, ushort, uchar)
506 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, short, uchar)
507 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, int, uchar)
508 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, uint, uchar)
509 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, float, uchar)
510 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, double, uchar)
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 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uchar, uchar)
527 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, char, char)
528 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, ushort, ushort)
529 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, short, short)
530 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, int, int)
531 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uint, uint)
532 
533 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_OP
534 
535 // binary operators (vec & scalar)
536 
537 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(op, input_type, scalar_type, output_type) \
538     __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, scalar_type s) \
539     { \
540         return VecTraits<output_type ## 1>::make(a.x op s); \
541     } \
542     __device__ __forceinline__ output_type ## 1 operator op(scalar_type s, const input_type ## 1 & b) \
543     { \
544         return VecTraits<output_type ## 1>::make(s op b.x); \
545     } \
546     __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, scalar_type s) \
547     { \
548         return VecTraits<output_type ## 2>::make(a.x op s, a.y op s); \
549     } \
550     __device__ __forceinline__ output_type ## 2 operator op(scalar_type s, const input_type ## 2 & b) \
551     { \
552         return VecTraits<output_type ## 2>::make(s op b.x, s op b.y); \
553     } \
554     __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, scalar_type s) \
555     { \
556         return VecTraits<output_type ## 3>::make(a.x op s, a.y op s, a.z op s); \
557     } \
558     __device__ __forceinline__ output_type ## 3 operator op(scalar_type s, const input_type ## 3 & b) \
559     { \
560         return VecTraits<output_type ## 3>::make(s op b.x, s op b.y, s op b.z); \
561     } \
562     __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, scalar_type s) \
563     { \
564         return VecTraits<output_type ## 4>::make(a.x op s, a.y op s, a.z op s, a.w op s); \
565     } \
566     __device__ __forceinline__ output_type ## 4 operator op(scalar_type s, const input_type ## 4 & b) \
567     { \
568         return VecTraits<output_type ## 4>::make(s op b.x, s op b.y, s op b.z, s op b.w); \
569     }
570 
571 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, int, int)
572 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, float, float)
573 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, double, double)
574 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, int, int)
575 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, float, float)
576 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, double, double)
577 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, int, int)
578 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, float, float)
579 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, double, double)
580 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, int, int)
581 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, float, float)
582 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, double, double)
583 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, int, int)
584 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, float, float)
585 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, double, double)
586 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, uint, uint)
587 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, float, float)
588 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, double, double)
589 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, float, float, float)
590 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, float, double, double)
591 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, double, double, double)
592 
593 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, int, int)
594 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, float, float)
595 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, double, double)
596 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, int, int)
597 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, float, float)
598 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, double, double)
599 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, int, int)
600 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, float, float)
601 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, double, double)
602 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, int, int)
603 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, float, float)
604 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, double, double)
605 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, int, int)
606 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, float, float)
607 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, double, double)
608 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, uint, uint)
609 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, float, float)
610 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, double, double)
611 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, float, float, float)
612 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, float, double, double)
613 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, double, double, double)
614 
615 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, int, int)
616 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, float, float)
617 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, double, double)
618 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, int, int)
619 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, float, float)
620 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, double, double)
621 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, int, int)
622 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, float, float)
623 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, double, double)
624 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, int, int)
625 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, float, float)
626 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, double, double)
627 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, int, int)
628 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, float, float)
629 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, double, double)
630 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, uint, uint)
631 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, float, float)
632 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, double, double)
633 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, float, float, float)
634 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, float, double, double)
635 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, double, double, double)
636 
637 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, int, int)
638 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, float, float)
639 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, double, double)
640 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, int, int)
641 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, float, float)
642 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, double, double)
643 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, int, int)
644 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, float, float)
645 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, double, double)
646 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, int, int)
647 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, float, float)
648 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, double, double)
649 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, int, int)
650 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, float, float)
651 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, double, double)
652 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, uint, uint)
653 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, float, float)
654 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, double, double)
655 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, float, float, float)
656 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, float, double, double)
657 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, double, double, double)
658 
659 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uchar, uchar, uchar)
660 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, char, char, uchar)
661 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, ushort, ushort, uchar)
662 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, short, short, uchar)
663 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, int, int, uchar)
664 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uint, uint, uchar)
665 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, float, float, uchar)
666 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, double, double, uchar)
667 
668 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uchar, uchar, uchar)
669 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, char, char, uchar)
670 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, ushort, ushort, uchar)
671 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, short, short, uchar)
672 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, int, int, uchar)
673 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uint, uint, uchar)
674 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, float, float, uchar)
675 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, double, double, uchar)
676 
677 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, uchar, uchar, uchar)
678 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, char, char, uchar)
679 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, ushort, ushort, uchar)
680 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, short, short, uchar)
681 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, int, int, uchar)
682 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, uint, uint, uchar)
683 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, float, float, uchar)
684 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, double, double, uchar)
685 
686 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, uchar, uchar, uchar)
687 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, char, char, uchar)
688 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, ushort, ushort, uchar)
689 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, short, short, uchar)
690 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, int, int, uchar)
691 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, uint, uint, uchar)
692 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, float, float, uchar)
693 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, double, double, uchar)
694 
695 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, uchar, uchar, uchar)
696 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, char, char, uchar)
697 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, ushort, ushort, uchar)
698 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, short, short, uchar)
699 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, int, int, uchar)
700 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, uint, uint, uchar)
701 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, float, float, uchar)
702 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, double, double, uchar)
703 
704 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, uchar, uchar, uchar)
705 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, char, char, uchar)
706 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, ushort, ushort, uchar)
707 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, short, short, uchar)
708 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, int, int, uchar)
709 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, uint, uint, uchar)
710 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, float, float, uchar)
711 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, double, double, uchar)
712 
713 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, uchar, uchar, uchar)
714 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, char, char, uchar)
715 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, ushort, ushort, uchar)
716 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, short, short, uchar)
717 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, int, int, uchar)
718 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, uint, uint, uchar)
719 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, float, float, uchar)
720 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, double, double, uchar)
721 
722 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, uchar, uchar, uchar)
723 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, char, char, uchar)
724 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, ushort, ushort, uchar)
725 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, short, short, uchar)
726 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, int, int, uchar)
727 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, uint, uint, uchar)
728 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, float, float, uchar)
729 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, double, double, uchar)
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 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uchar, uchar, uchar)
746 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, char, char, char)
747 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, ushort, ushort, ushort)
748 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, short, short, short)
749 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, int, int, int)
750 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uint, uint, uint)
751 
752 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP
753 
754 // binary function (vec & vec)
755 
756 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(func_name, func, input_type, output_type) \
757     __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, const input_type ## 1 & b) \
758     { \
759         return VecTraits<output_type ## 1>::make(func (a.x, b.x)); \
760     } \
761     __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, const input_type ## 2 & b) \
762     { \
763         return VecTraits<output_type ## 2>::make(func (a.x, b.x), func (a.y, b.y)); \
764     } \
765     __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, const input_type ## 3 & b) \
766     { \
767         return VecTraits<output_type ## 3>::make(func (a.x, b.x), func (a.y, b.y), func (a.z, b.z)); \
768     } \
769     __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, const input_type ## 4 & b) \
770     { \
771         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)); \
772     }
773 
774 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, uchar, uchar)
775 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, char, char)
776 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, ushort, ushort)
777 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, short, short)
778 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, uint, uint)
779 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, int, int)
780 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::fmaxf, float, float)
781 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::fmax, double, double)
782 
783 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, uchar, uchar)
784 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, char, char)
785 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, ushort, ushort)
786 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, short, short)
787 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, uint, uint)
788 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, int, int)
789 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::fminf, float, float)
790 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::fmin, double, double)
791 
792 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, uchar, float)
793 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, char, float)
794 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, ushort, float)
795 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, short, float)
796 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, uint, float)
797 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, int, float)
798 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, float, float)
799 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypot, double, double)
800 
801 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, uchar, float)
802 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, char, float)
803 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, ushort, float)
804 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, short, float)
805 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, uint, float)
806 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, int, float)
807 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, float, float)
808 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2, double, double)
809 
810 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC
811 
812 // binary function (vec & scalar)
813 
814 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(func_name, func, input_type, scalar_type, output_type) \
815     __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, scalar_type s) \
816     { \
817         return VecTraits<output_type ## 1>::make(func ((output_type) a.x, (output_type) s)); \
818     } \
819     __device__ __forceinline__ output_type ## 1 func_name(scalar_type s, const input_type ## 1 & b) \
820     { \
821         return VecTraits<output_type ## 1>::make(func ((output_type) s, (output_type) b.x)); \
822     } \
823     __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, scalar_type s) \
824     { \
825         return VecTraits<output_type ## 2>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s)); \
826     } \
827     __device__ __forceinline__ output_type ## 2 func_name(scalar_type s, const input_type ## 2 & b) \
828     { \
829         return VecTraits<output_type ## 2>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y)); \
830     } \
831     __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, scalar_type s) \
832     { \
833         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)); \
834     } \
835     __device__ __forceinline__ output_type ## 3 func_name(scalar_type s, const input_type ## 3 & b) \
836     { \
837         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)); \
838     } \
839     __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, scalar_type s) \
840     { \
841         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)); \
842     } \
843     __device__ __forceinline__ output_type ## 4 func_name(scalar_type s, const input_type ## 4 & b) \
844     { \
845         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)); \
846     }
847 
848 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, uchar, uchar, uchar)
849 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, uchar, float, float)
850 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, uchar, double, double)
851 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, char, char, char)
852 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, char, float, float)
853 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, char, double, double)
854 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, ushort, ushort, ushort)
855 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, ushort, float, float)
856 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, ushort, double, double)
857 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, short, short, short)
858 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, short, float, float)
859 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, short, double, double)
860 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, uint, uint, uint)
861 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, uint, float, float)
862 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, uint, double, double)
863 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, int, int, int)
864 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, int, float, float)
865 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, int, double, double)
866 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, float, float, float)
867 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, float, double, double)
868 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, double, double, double)
869 
870 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, uchar, uchar, uchar)
871 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, uchar, float, float)
872 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, uchar, double, double)
873 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, char, char, char)
874 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, char, float, float)
875 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, char, double, double)
876 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, ushort, ushort, ushort)
877 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, ushort, float, float)
878 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, ushort, double, double)
879 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, short, short, short)
880 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, short, float, float)
881 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, short, double, double)
882 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, uint, uint, uint)
883 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, uint, float, float)
884 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, uint, double, double)
885 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, int, int, int)
886 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, int, float, float)
887 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, int, double, double)
888 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, float, float, float)
889 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, float, double, double)
890 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, double, double, double)
891 
892 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, uchar, float, float)
893 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, uchar, double, double)
894 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, char, float, float)
895 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, char, double, double)
896 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, ushort, float, float)
897 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, ushort, double, double)
898 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, short, float, float)
899 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, short, double, double)
900 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, uint, float, float)
901 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, uint, double, double)
902 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, int, float, float)
903 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, int, double, double)
904 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, float, float, float)
905 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, float, double, double)
906 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, double, double, double)
907 
908 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, uchar, float, float)
909 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, uchar, double, double)
910 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, char, float, float)
911 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, char, double, double)
912 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, ushort, float, float)
913 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, ushort, double, double)
914 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, short, float, float)
915 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, short, double, double)
916 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, uint, float, float)
917 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, uint, double, double)
918 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, int, float, float)
919 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, int, double, double)
920 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, float, float, float)
921 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, float, double, double)
922 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, double, double, double)
923 
924 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC
925 
926 }}} // namespace cv { namespace cuda { namespace device
927 
928 //! @endcond
929 
930 #endif // __OPENCV_CUDA_VECMATH_HPP__
931