43 #ifndef __OPENCV_GPU_VECMATH_HPP__
44 #define __OPENCV_GPU_VECMATH_HPP__
49 namespace cv {
namespace gpu {
namespace device
54 namespace vec_math_detail
59 template <
typename VecS>
static __device__ __forceinline__ VecD
cast(
const VecS&
v)
67 template <
typename VecS>
static __device__ __forceinline__ VecD
cast(
const VecS&
v)
75 template <
typename VecS>
static __device__ __forceinline__ VecD
cast(
const VecS&
v)
83 template <
typename VecS>
static __device__ __forceinline__ VecD
cast(
const VecS&
v)
86 return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y), saturate_cast<D>(v.z), saturate_cast<D>(v.w));
90 template <
typename VecD,
typename VecS>
static __device__ __forceinline__ VecD saturate_cast_helper(
const VecS&
v)
92 return SatCastHelper<VecTraits<VecD>::cn, VecD>::cast(v);
96 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uchar1&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
97 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const char1&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
98 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const ushort1&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
99 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const short1&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
100 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uint1&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
101 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const int1&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
102 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const float1&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
103 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const double1&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
105 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uchar2&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
106 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const char2&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
107 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const ushort2&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
108 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const short2&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
109 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uint2&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
110 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const int2&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
111 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const float2&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
112 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const double2&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
114 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uchar3&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
115 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const char3&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
116 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const ushort3&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
117 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const short3&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
118 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uint3&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
119 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const int3&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
120 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const float3&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
121 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const double3&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
123 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uchar4&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
124 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const char4&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
125 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const ushort4&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
126 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const short4&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
127 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const uint4&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
128 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const int4&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
129 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const float4&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
130 template<
typename T>
static __device__ __forceinline__
T saturate_cast(
const double4&
v) {
return vec_math_detail::saturate_cast_helper<T>(v);}
134 #define CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(op, input_type, output_type) \
135 __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a) \
137 return VecTraits<output_type ## 1>::make(op (a.x)); \
139 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a) \
141 return VecTraits<output_type ## 2>::make(op (a.x), op (a.y)); \
143 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a) \
145 return VecTraits<output_type ## 3>::make(op (a.x), op (a.y), op (a.z)); \
147 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a) \
149 return VecTraits<output_type ## 4>::make(op (a.x), op (a.y), op (a.z), op (a.w)); \
155 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-,
float,
float)
156 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-,
double,
double)
158 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!,
uchar, uchar)
159 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!,
char, uchar)
160 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!,
ushort, uchar)
161 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!,
short, uchar)
162 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!,
int, uchar)
163 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!,
uint, uchar)
164 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!,
float, uchar)
165 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!,
double, uchar)
167 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, uchar, uchar)
168 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~,
char,
char)
169 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, ushort, ushort)
170 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~,
short,
short)
171 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~,
int,
int)
172 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, uint, uint)
174 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_OP
178 #define CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(func_name, func, input_type, output_type) \
179 __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a) \
181 return VecTraits<output_type ## 1>::make(func (a.x)); \
183 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a) \
185 return VecTraits<output_type ## 2>::make(func (a.x), func (a.y)); \
187 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a) \
189 return VecTraits<output_type ## 3>::make(func (a.x), func (a.y), func (a.z)); \
191 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a) \
193 return VecTraits<output_type ## 4>::make(func (a.x), func (a.y), func (a.z), func (a.w)); \
376 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC
380 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(op, input_type, output_type) \
381 __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, const input_type ## 1 & b) \
383 return VecTraits<output_type ## 1>::make(a.x op b.x); \
385 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, const input_type ## 2 & b) \
387 return VecTraits<output_type ## 2>::make(a.x op b.x, a.y op b.y); \
389 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, const input_type ## 3 & b) \
391 return VecTraits<output_type ## 3>::make(a.x op b.x, a.y op b.y, a.z op b.z); \
393 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, const input_type ## 4 & b) \
395 return VecTraits<output_type ## 4>::make(a.x op b.x, a.y op b.y, a.z op b.z, a.w op b.w); \
400 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, ushort,
int)
401 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+,
short,
int)
402 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+,
int,
int)
403 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, uint, uint)
404 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+,
float,
float)
405 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+,
double,
double)
407 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uchar,
int)
408 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-,
char,
int)
409 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, ushort,
int)
410 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-,
short,
int)
411 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-,
int,
int)
412 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uint, uint)
413 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-,
float,
float)
414 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-,
double,
double)
416 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, uchar,
int)
417 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*,
char,
int)
418 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, ushort,
int)
419 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*,
short,
int)
420 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*,
int,
int)
421 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, uint, uint)
422 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*,
float,
float)
423 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*,
double,
double)
425 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, uchar,
int)
426 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/,
char,
int)
427 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, ushort,
int)
428 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/,
short,
int)
429 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/,
int,
int)
430 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, uint, uint)
431 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/,
float,
float)
432 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/,
double,
double)
434 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uchar, uchar)
435 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==,
char, uchar)
436 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, ushort, uchar)
437 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==,
short, uchar)
438 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==,
int, uchar)
439 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uint, uchar)
440 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==,
float, uchar)
441 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==,
double, uchar)
443 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uchar, uchar)
444 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=,
char, uchar)
445 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, ushort, uchar)
446 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=,
short, uchar)
447 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=,
int, uchar)
448 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uint, uchar)
449 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=,
float, uchar)
450 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=,
double, uchar)
452 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, uchar, uchar)
453 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>,
char, uchar)
454 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, ushort, uchar)
455 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>,
short, uchar)
456 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>,
int, uchar)
457 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, uint, uchar)
458 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>,
float, uchar)
459 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>,
double, uchar)
461 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, uchar, uchar)
462 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<,
char, uchar)
463 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, ushort, uchar)
464 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<,
short, uchar)
465 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<,
int, uchar)
466 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, uint, uchar)
467 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<,
float, uchar)
468 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<,
double, uchar)
470 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, uchar, uchar)
471 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=,
char, uchar)
472 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, ushort, uchar)
473 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=,
short, uchar)
474 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=,
int, uchar)
475 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, uint, uchar)
476 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=,
float, uchar)
477 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=,
double, uchar)
479 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, uchar, uchar)
480 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=,
char, uchar)
481 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, ushort, uchar)
482 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=,
short, uchar)
483 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=,
int, uchar)
484 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, uint, uchar)
485 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=,
float, uchar)
486 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=,
double, uchar)
488 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, uchar, uchar)
489 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&,
char, uchar)
490 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, ushort, uchar)
491 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&,
short, uchar)
492 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&,
int, uchar)
493 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, uint, uchar)
494 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&,
float, uchar)
495 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&,
double, uchar)
497 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, uchar, uchar)
498 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||,
char, uchar)
499 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, ushort, uchar)
500 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||,
short, uchar)
501 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||,
int, uchar)
502 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, uint, uchar)
503 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||,
float, uchar)
504 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||,
double, uchar)
506 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, uchar, uchar)
507 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&,
char,
char)
508 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, ushort, ushort)
509 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&,
short,
short)
510 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&,
int,
int)
511 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, uint, uint)
513 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, uchar, uchar)
514 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|,
char,
char)
515 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, ushort, ushort)
516 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|,
short,
short)
517 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|,
int,
int)
518 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, uint, uint)
520 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uchar, uchar)
521 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^,
char,
char)
522 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, ushort, ushort)
523 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^,
short,
short)
524 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^,
int,
int)
525 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uint, uint)
527 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_OP
531 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(op, input_type, scalar_type, output_type) \
532 __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, scalar_type s) \
534 return VecTraits<output_type ## 1>::make(a.x op s); \
536 __device__ __forceinline__ output_type ## 1 operator op(scalar_type s, const input_type ## 1 & b) \
538 return VecTraits<output_type ## 1>::make(s op b.x); \
540 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, scalar_type s) \
542 return VecTraits<output_type ## 2>::make(a.x op s, a.y op s); \
544 __device__ __forceinline__ output_type ## 2 operator op(scalar_type s, const input_type ## 2 & b) \
546 return VecTraits<output_type ## 2>::make(s op b.x, s op b.y); \
548 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, scalar_type s) \
550 return VecTraits<output_type ## 3>::make(a.x op s, a.y op s, a.z op s); \
552 __device__ __forceinline__ output_type ## 3 operator op(scalar_type s, const input_type ## 3 & b) \
554 return VecTraits<output_type ## 3>::make(s op b.x, s op b.y, s op b.z); \
556 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, scalar_type s) \
558 return VecTraits<output_type ## 4>::make(a.x op s, a.y op s, a.z op s, a.w op s); \
560 __device__ __forceinline__ output_type ## 4 operator op(scalar_type s, const input_type ## 4 & b) \
562 return VecTraits<output_type ## 4>::make(s op b.x, s op b.y, s op b.z, s op b.w); \
567 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar,
double,
double)
568 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
char,
int,
int)
569 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
char,
float,
float)
570 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
char,
double,
double)
571 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort,
int,
int)
572 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort,
float,
float)
573 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort,
double,
double)
574 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
short,
int,
int)
575 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
short,
float,
float)
576 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
short,
double,
double)
577 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
int,
int,
int)
578 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
int,
float,
float)
579 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
int,
double,
double)
580 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, uint, uint)
581 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint,
float,
float)
582 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint,
double,
double)
583 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
float,
float,
float)
584 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
float,
double,
double)
585 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+,
double,
double,
double)
587 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar,
int,
int)
588 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar,
float,
float)
589 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar,
double,
double)
590 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
char,
int,
int)
591 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
char,
float,
float)
592 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
char,
double,
double)
593 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort,
int,
int)
594 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort,
float,
float)
595 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort,
double,
double)
596 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
short,
int,
int)
597 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
short,
float,
float)
598 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
short,
double,
double)
599 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
int,
int,
int)
600 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
int,
float,
float)
601 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
int,
double,
double)
602 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, uint, uint)
603 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint,
float,
float)
604 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint,
double,
double)
605 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
float,
float,
float)
606 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
float,
double,
double)
607 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-,
double,
double,
double)
609 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar,
int,
int)
610 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar,
float,
float)
611 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar,
double,
double)
612 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
char,
int,
int)
613 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
char,
float,
float)
614 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
char,
double,
double)
615 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort,
int,
int)
616 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort,
float,
float)
617 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort,
double,
double)
618 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
short,
int,
int)
619 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
short,
float,
float)
620 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
short,
double,
double)
621 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
int,
int,
int)
622 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
int,
float,
float)
623 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
int,
double,
double)
624 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, uint, uint)
625 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint,
float,
float)
626 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint,
double,
double)
627 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
float,
float,
float)
628 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
float,
double,
double)
629 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*,
double,
double,
double)
631 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar,
int,
int)
632 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar,
float,
float)
633 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar,
double,
double)
634 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
char,
int,
int)
635 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
char,
float,
float)
636 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
char,
double,
double)
637 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort,
int,
int)
638 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort,
float,
float)
639 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort,
double,
double)
640 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
short,
int,
int)
641 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
short,
float,
float)
642 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
short,
double,
double)
643 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
int,
int,
int)
644 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
int,
float,
float)
645 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
int,
double,
double)
646 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, uint, uint)
647 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint,
float,
float)
648 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint,
double,
double)
649 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
float,
float,
float)
650 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
float,
double,
double)
651 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/,
double,
double,
double)
653 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uchar, uchar, uchar)
654 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==,
char,
char, uchar)
655 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, ushort, ushort, uchar)
656 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==,
short,
short, uchar)
657 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==,
int,
int, uchar)
658 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uint, uint, uchar)
659 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==,
float,
float, uchar)
660 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==,
double,
double, uchar)
662 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uchar, uchar, uchar)
663 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=,
char,
char, uchar)
664 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, ushort, ushort, uchar)
665 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=,
short,
short, uchar)
666 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=,
int,
int, uchar)
667 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uint, uint, uchar)
668 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=,
float,
float, uchar)
669 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=,
double,
double, uchar)
671 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, uchar, uchar, uchar)
672 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>,
char,
char, uchar)
673 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, ushort, ushort, uchar)
674 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>,
short,
short, uchar)
675 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>,
int,
int, uchar)
676 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, uint, uint, uchar)
677 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>,
float,
float, uchar)
678 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>,
double,
double, uchar)
680 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, uchar, uchar, uchar)
681 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<,
char,
char, uchar)
682 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, ushort, ushort, uchar)
683 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<,
short,
short, uchar)
684 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<,
int,
int, uchar)
685 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, uint, uint, uchar)
686 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<,
float,
float, uchar)
687 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<,
double,
double, uchar)
689 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, uchar, uchar, uchar)
690 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=,
char,
char, uchar)
691 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, ushort, ushort, uchar)
692 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=,
short,
short, uchar)
693 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=,
int,
int, uchar)
694 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, uint, uint, uchar)
695 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=,
float,
float, uchar)
696 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=,
double,
double, uchar)
698 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, uchar, uchar, uchar)
699 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=,
char,
char, uchar)
700 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, ushort, ushort, uchar)
701 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=,
short,
short, uchar)
702 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=,
int,
int, uchar)
703 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, uint, uint, uchar)
704 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=,
float,
float, uchar)
705 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=,
double,
double, uchar)
707 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, uchar, uchar, uchar)
708 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&,
char,
char, uchar)
709 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, ushort, ushort, uchar)
710 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&,
short,
short, uchar)
711 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&,
int,
int, uchar)
712 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, uint, uint, uchar)
713 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&,
float,
float, uchar)
714 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&,
double,
double, uchar)
716 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, uchar, uchar, uchar)
717 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||,
char,
char, uchar)
718 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, ushort, ushort, uchar)
719 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||,
short,
short, uchar)
720 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||,
int,
int, uchar)
721 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, uint, uint, uchar)
722 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||,
float,
float, uchar)
723 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||,
double,
double, uchar)
725 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, uchar, uchar, uchar)
726 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&,
char,
char,
char)
727 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, ushort, ushort, ushort)
728 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&,
short,
short,
short)
729 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&,
int,
int,
int)
730 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, uint, uint, uint)
732 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, uchar, uchar, uchar)
733 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|,
char,
char,
char)
734 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, ushort, ushort, ushort)
735 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|,
short,
short,
short)
736 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|,
int,
int,
int)
737 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, uint, uint, uint)
739 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uchar, uchar, uchar)
740 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^,
char,
char,
char)
741 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, ushort, ushort, ushort)
742 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^,
short,
short,
short)
743 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^,
int,
int,
int)
744 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uint, uint, uint)
746 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP
750 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(func_name, func, input_type, output_type) \
751 __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, const input_type ## 1 & b) \
753 return VecTraits<output_type ## 1>::make(func (a.x, b.x)); \
755 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, const input_type ## 2 & b) \
757 return VecTraits<output_type ## 2>::make(func (a.x, b.x), func (a.y, b.y)); \
759 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, const input_type ## 3 & b) \
761 return VecTraits<output_type ## 3>::make(func (a.x, b.x), func (a.y, b.y), func (a.z, b.z)); \
763 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, const input_type ## 4 & b) \
765 return VecTraits<output_type ## 4>::make(func (a.x, b.x), func (a.y, b.y), func (a.z, b.z), func (a.w, b.w)); \
804 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC
808 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(func_name, func, input_type, scalar_type, output_type) \
809 __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, scalar_type s) \
811 return VecTraits<output_type ## 1>::make(func ((output_type) a.x, (output_type) s)); \
813 __device__ __forceinline__ output_type ## 1 func_name(scalar_type s, const input_type ## 1 & b) \
815 return VecTraits<output_type ## 1>::make(func ((output_type) s, (output_type) b.x)); \
817 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, scalar_type s) \
819 return VecTraits<output_type ## 2>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s)); \
821 __device__ __forceinline__ output_type ## 2 func_name(scalar_type s, const input_type ## 2 & b) \
823 return VecTraits<output_type ## 2>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y)); \
825 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, scalar_type s) \
827 return VecTraits<output_type ## 3>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s), func ((output_type) a.z, (output_type) s)); \
829 __device__ __forceinline__ output_type ## 3 func_name(scalar_type s, const input_type ## 3 & b) \
831 return VecTraits<output_type ## 3>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y), func ((output_type) s, (output_type) b.z)); \
833 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, scalar_type s) \
835 return VecTraits<output_type ## 4>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s), func ((output_type) a.z, (output_type) s), func ((output_type) a.w, (output_type) s)); \
837 __device__ __forceinline__ output_type ## 4 func_name(scalar_type s, const input_type ## 4 & b) \
839 return VecTraits<output_type ## 4>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y), func ((output_type) s, (output_type) b.z), func ((output_type) s, (output_type) b.w)); \
918 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC
922 #endif // __OPENCV_GPU_VECMATH_HPP__
short float uchar uchar uchar uchar uchar ushort int uchar ushort int float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float float int CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, char, int) CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+
CV_EXPORTS void sqrt(const GpuMat &src, GpuMat &dst, Stream &stream=Stream::Null())
CV_EXPORTS void exp(const GpuMat &a, GpuMat &b, Stream &stream=Stream::Null())
static __device__ __forceinline__ VecD cast(const VecS &v)
Definition: vec_math.hpp:83
static __device__ __forceinline__ VecD cast(const VecS &v)
Definition: vec_math.hpp:75
CV_EXPORTS void max(const GpuMat &src1, const GpuMat &src2, GpuMat &dst, Stream &stream=Stream::Null())
computes per-element maximum of two arrays (dst = max(src1, src2))
CV_EXPORTS void min(const GpuMat &src1, const GpuMat &src2, GpuMat &dst, Stream &stream=Stream::Null())
computes per-element minimum of two arrays (dst = min(src1, src2))
CV_EXPORTS void abs(const GpuMat &src, GpuMat &dst, Stream &stream=Stream::Null())
CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max,::max, uchar, uchar, uchar) CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max
const CvMat const CvMat const CvMat CvMat CvMat CvMat CvMat CvSize CvMat CvMat * T
Definition: calib3d.hpp:270
CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, int, int) CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+
Definition: vec_math.hpp:56
short float uchar uchar uchar uchar uchar ushort int uchar CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs,::abs, char, char) CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs
unsigned char uchar
Definition: common.hpp:100
static __device__ __forceinline__ VecD cast(const VecS &v)
Definition: vec_math.hpp:67
CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, char, char) CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-
unsigned int uint
Definition: common.hpp:104
unsigned short ushort
Definition: common.hpp:101
__device__ __forceinline__ _Tp saturate_cast(uchar v)
Definition: saturate_cast.hpp:50
CV_EXPORTS void log(const GpuMat &a, GpuMat &b, Stream &stream=Stream::Null())
static __device__ __forceinline__ VecD cast(const VecS &v)
Definition: vec_math.hpp:59
Definition: vec_traits.hpp:160
CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max,::max, uchar, uchar) CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max