43#ifndef OPENCV_CUDA_VECMATH_HPP
44#define OPENCV_CUDA_VECMATH_HPP
46#include "vec_traits.hpp"
47#include "saturate_cast.hpp"
55namespace cv {
namespace cuda {
namespace device
60namespace vec_math_detail
62 template <
int cn,
typename VecD>
struct SatCastHelper;
63 template <
typename VecD>
struct SatCastHelper<1, VecD>
65 template <
typename VecS>
static __device__ __forceinline__ VecD cast(
const VecS& v)
67 typedef typename VecTraits<VecD>::elem_type D;
68 return VecTraits<VecD>::make(saturate_cast<D>(v.x));
71 template <
typename VecD>
struct SatCastHelper<2, VecD>
73 template <
typename VecS>
static __device__ __forceinline__ VecD cast(
const VecS& v)
75 typedef typename VecTraits<VecD>::elem_type D;
76 return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y));
79 template <
typename VecD>
struct SatCastHelper<3, VecD>
81 template <
typename VecS>
static __device__ __forceinline__ VecD cast(
const VecS& v)
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));
87 template <
typename VecD>
struct SatCastHelper<4, VecD>
89 template <
typename VecS>
static __device__ __forceinline__ VecD cast(
const VecS& v)
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));
96 template <
typename VecD,
typename VecS>
static __device__ __forceinline__ VecD saturate_cast_helper(
const VecS& v)
98 return SatCastHelper<VecTraits<VecD>::cn, VecD>::cast(v);
102template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uchar1& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
103template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const char1& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
104template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const ushort1& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
105template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const short1& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
106template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uint1& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
107template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const int1& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
108template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const float1& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
109template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const double1& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
111template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uchar2& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
112template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const char2& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
113template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const ushort2& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
114template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const short2& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
115template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uint2& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
116template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const int2& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
117template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const float2& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
118template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const double2& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
120template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uchar3& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
121template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const char3& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
122template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const ushort3& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
123template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const short3& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
124template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uint3& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
125template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const int3& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
126template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const float3& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
127template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const double3& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
129template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uchar4& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
130template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const char4& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
131template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const ushort4& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
132template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const short4& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
133template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uint4& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
134template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const int4& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
135template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const float4& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
136template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const double4& v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
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) \
143 return VecTraits<output_type ## 1>::make(op (a.x)); \
145 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a) \
147 return VecTraits<output_type ## 2>::make(op (a.x), op (a.y)); \
149 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a) \
151 return VecTraits<output_type ## 3>::make(op (a.x), op (a.y), op (a.z)); \
153 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a) \
155 return VecTraits<output_type ## 4>::make(op (a.x), op (a.y), op (a.z), op (a.w)); \
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)
164CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!,
uchar,
uchar)
165CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!,
char,
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)
173CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~,
uchar,
uchar)
174CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~,
char,
char)
176CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~,
short,
short)
177CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~,
int,
int)
178CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~,
uint,
uint)
180#undef CV_CUDEV_IMPLEMENT_VEC_UNARY_OP
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) \
187 return VecTraits<output_type ## 1>::make(func (a.x)); \
189 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a) \
191 return VecTraits<output_type ## 2>::make(func (a.x), func (a.y)); \
193 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a) \
195 return VecTraits<output_type ## 3>::make(func (a.x), func (a.y), func (a.z)); \
197 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a) \
199 return VecTraits<output_type ## 4>::make(func (a.x), func (a.y), func (a.z), func (a.w)); \
202CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::fabsf,
float,
float)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
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)
375#undef CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC
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) \
382 return VecTraits<output_type ## 1>::make(a.x op b.x); \
384 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, const input_type ## 2 & b) \
386 return VecTraits<output_type ## 2>::make(a.x op b.x, a.y op b.y); \
388 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, const input_type ## 3 & b) \
390 return VecTraits<output_type ## 3>::make(a.x op b.x, a.y op b.y, a.z op b.z); \
392 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, const input_type ## 4 & b) \
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); \
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)
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)
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)
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)
433CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==,
uchar,
uchar)
434CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==,
char,
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)
442CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=,
uchar,
uchar)
443CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=,
char,
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)
451CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>,
uchar,
uchar)
452CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>,
char,
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)
460CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<,
uchar,
uchar)
461CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<,
char,
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)
469CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=,
uchar,
uchar)
470CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=,
char,
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)
478CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=,
uchar,
uchar)
479CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=,
char,
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)
487CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&,
uchar,
uchar)
488CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&,
char,
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)
496CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||,
uchar,
uchar)
497CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||,
char,
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)
505CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&,
uchar,
uchar)
506CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&,
char,
char)
508CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&,
short,
short)
509CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&,
int,
int)
510CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&,
uint,
uint)
512CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|,
uchar,
uchar)
513CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|,
char,
char)
515CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|,
short,
short)
516CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|,
int,
int)
517CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|,
uint,
uint)
519CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^,
uchar,
uchar)
520CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^,
char,
char)
522CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^,
short,
short)
523CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^,
int,
int)
524CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^,
uint,
uint)
526#undef CV_CUDEV_IMPLEMENT_VEC_BINARY_OP
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) \
533 return VecTraits<output_type ## 1>::make(a.x op s); \
535 __device__ __forceinline__ output_type ## 1 operator op(scalar_type s, const input_type ## 1 & b) \
537 return VecTraits<output_type ## 1>::make(s op b.x); \
539 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, scalar_type s) \
541 return VecTraits<output_type ## 2>::make(a.x op s, a.y op s); \
543 __device__ __forceinline__ output_type ## 2 operator op(scalar_type s, const input_type ## 2 & b) \
545 return VecTraits<output_type ## 2>::make(s op b.x, s op b.y); \
547 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, scalar_type s) \
549 return VecTraits<output_type ## 3>::make(a.x op s, a.y op s, a.z op s); \
551 __device__ __forceinline__ output_type ## 3 operator op(scalar_type s, const input_type ## 3 & b) \
553 return VecTraits<output_type ## 3>::make(s op b.x, s op b.y, s op b.z); \
555 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, scalar_type s) \
557 return VecTraits<output_type ## 4>::make(a.x op s, a.y op s, a.z op s, a.w op s); \
559 __device__ __forceinline__ output_type ## 4 operator op(scalar_type s, const input_type ## 4 & b) \
561 return VecTraits<output_type ## 4>::make(s op b.x, s op b.y, s op b.z, s op b.w); \
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)
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)
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)
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)
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)
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)
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)
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)
653CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==,
char,
char,
uchar)
655CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==,
short,
short,
uchar)
656CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==,
int,
int,
uchar)
658CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==,
float,
float,
uchar)
659CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==,
double,
double,
uchar)
662CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=,
char,
char,
uchar)
664CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=,
short,
short,
uchar)
665CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=,
int,
int,
uchar)
667CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=,
float,
float,
uchar)
668CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=,
double,
double,
uchar)
671CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>,
char,
char,
uchar)
673CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>,
short,
short,
uchar)
674CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>,
int,
int,
uchar)
676CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>,
float,
float,
uchar)
677CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>,
double,
double,
uchar)
680CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<,
char,
char,
uchar)
682CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<,
short,
short,
uchar)
683CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<,
int,
int,
uchar)
685CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<,
float,
float,
uchar)
686CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<,
double,
double,
uchar)
689CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=,
char,
char,
uchar)
691CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=,
short,
short,
uchar)
692CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=,
int,
int,
uchar)
694CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=,
float,
float,
uchar)
695CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=,
double,
double,
uchar)
698CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=,
char,
char,
uchar)
700CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=,
short,
short,
uchar)
701CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=,
int,
int,
uchar)
703CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=,
float,
float,
uchar)
704CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=,
double,
double,
uchar)
707CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&,
char,
char,
uchar)
709CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&,
short,
short,
uchar)
710CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&,
int,
int,
uchar)
712CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&,
float,
float,
uchar)
713CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&,
double,
double,
uchar)
716CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||,
char,
char,
uchar)
718CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||,
short,
short,
uchar)
719CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||,
int,
int,
uchar)
721CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||,
float,
float,
uchar)
722CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||,
double,
double,
uchar)
725CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&,
char,
char,
char)
727CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&,
short,
short,
short)
728CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&,
int,
int,
int)
732CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|,
char,
char,
char)
734CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|,
short,
short,
short)
735CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|,
int,
int,
int)
739CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^,
char,
char,
char)
741CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^,
short,
short,
short)
742CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^,
int,
int,
int)
745#undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP
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) \
752 return VecTraits<output_type ## 1>::make(func (a.x, b.x)); \
754 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, const input_type ## 2 & b) \
756 return VecTraits<output_type ## 2>::make(func (a.x, b.x), func (a.y, b.y)); \
758 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, const input_type ## 3 & b) \
760 return VecTraits<output_type ## 3>::make(func (a.x, b.x), func (a.y, b.y), func (a.z, b.z)); \
762 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, const input_type ## 4 & b) \
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)); \
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)
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)
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)
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)
803#undef CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC
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) \
810 return VecTraits<output_type ## 1>::make(func ((output_type) a.x, (output_type) s)); \
812 __device__ __forceinline__ output_type ## 1 func_name(scalar_type s, const input_type ## 1 & b) \
814 return VecTraits<output_type ## 1>::make(func ((output_type) s, (output_type) b.x)); \
816 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, scalar_type s) \
818 return VecTraits<output_type ## 2>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s)); \
820 __device__ __forceinline__ output_type ## 2 func_name(scalar_type s, const input_type ## 2 & b) \
822 return VecTraits<output_type ## 2>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y)); \
824 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, scalar_type s) \
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)); \
828 __device__ __forceinline__ output_type ## 3 func_name(scalar_type s, const input_type ## 3 & b) \
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)); \
832 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, scalar_type s) \
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)); \
836 __device__ __forceinline__ output_type ## 4 func_name(scalar_type s, const input_type ## 4 & b) \
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)); \
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)
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)
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)
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)
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)
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)
917#undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC
InputArrayOfArrays InputArrayOfArrays InputOutputArray InputOutputArray InputOutputArray InputOutputArray Size InputOutputArray InputOutputArray T
Definition calib3d.hpp:1867
uint32_t uint
Definition interface.h:42
unsigned char uchar
Definition interface.h:51
unsigned short ushort
Definition interface.h:52
"black box" representation of the file storage associated with a file on disk.
Definition calib3d.hpp:441