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