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