Renesas GR-PEACH OpenCV Development
/
gr-peach-opencv-project-sd-card_update
Important changes to repositories hosted on mbed.com
Mbed hosted mercurial repositories are deprecated and are due to be permanently deleted in July 2026.
To keep a copy of this software download the repository Zip archive or clone locally using Mercurial.
It is also possible to export all your personal repositories from the account settings page.
Fork of
gr-peach-opencv-project-sd-card
by 
the do
Embed:
(wiki syntax)
Show/hide line numbers
vec_math.hpp
Go to the documentation of this file.
00001 /*M/////////////////////////////////////////////////////////////////////////////////////// 00002 // 00003 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. 00004 // 00005 // By downloading, copying, installing or using the software you agree to this license. 00006 // If you do not agree to this license, do not download, install, 00007 // copy or use the software. 00008 // 00009 // 00010 // License Agreement 00011 // For Open Source Computer Vision Library 00012 // 00013 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved. 00014 // Copyright (C) 2009, Willow Garage Inc., all rights reserved. 00015 // Third party copyrights are property of their respective owners. 00016 // 00017 // Redistribution and use in source and binary forms, with or without modification, 00018 // are permitted provided that the following conditions are met: 00019 // 00020 // * Redistribution's of source code must retain the above copyright notice, 00021 // this list of conditions and the following disclaimer. 00022 // 00023 // * Redistribution's in binary form must reproduce the above copyright notice, 00024 // this list of conditions and the following disclaimer in the documentation 00025 // and/or other materials provided with the distribution. 00026 // 00027 // * The name of the copyright holders may not be used to endorse or promote products 00028 // derived from this software without specific prior written permission. 00029 // 00030 // This software is provided by the copyright holders and contributors "as is" and 00031 // any express or implied warranties, including, but not limited to, the implied 00032 // warranties of merchantability and fitness for a particular purpose are disclaimed. 00033 // In no event shall the Intel Corporation or contributors be liable for any direct, 00034 // indirect, incidental, special, exemplary, or consequential damages 00035 // (including, but not limited to, procurement of substitute goods or services; 00036 // loss of use, data, or profits; or business interruption) however caused 00037 // and on any theory of liability, whether in contract, strict liability, 00038 // or tort (including negligence or otherwise) arising in any way out of 00039 // the use of this software, even if advised of the possibility of such damage. 00040 // 00041 //M*/ 00042 00043 #ifndef __OPENCV_CUDA_VECMATH_HPP__ 00044 #define __OPENCV_CUDA_VECMATH_HPP__ 00045 00046 #include "vec_traits.hpp " 00047 #include "saturate_cast.hpp " 00048 00049 /** @file 00050 * @deprecated Use @ref cudev instead. 00051 */ 00052 00053 //! @cond IGNORED 00054 00055 namespace cv { namespace cuda { namespace device 00056 { 00057 00058 // saturate_cast 00059 00060 namespace vec_math_detail 00061 { 00062 template <int cn, typename VecD> struct SatCastHelper; 00063 template <typename VecD> struct SatCastHelper<1, VecD> 00064 { 00065 template <typename VecS> static __device__ __forceinline__ VecD cast(const VecS& v) 00066 { 00067 typedef typename VecTraits<VecD>::elem_type D; 00068 return VecTraits<VecD>::make(saturate_cast<D>(v.x)); 00069 } 00070 }; 00071 template <typename VecD> struct SatCastHelper<2, VecD> 00072 { 00073 template <typename VecS> static __device__ __forceinline__ VecD cast(const VecS& v) 00074 { 00075 typedef typename VecTraits<VecD>::elem_type D; 00076 return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y)); 00077 } 00078 }; 00079 template <typename VecD> struct SatCastHelper<3, VecD> 00080 { 00081 template <typename VecS> static __device__ __forceinline__ VecD cast(const VecS& v) 00082 { 00083 typedef typename VecTraits<VecD>::elem_type D; 00084 return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y), saturate_cast<D>(v.z)); 00085 } 00086 }; 00087 template <typename VecD> struct SatCastHelper<4, VecD> 00088 { 00089 template <typename VecS> static __device__ __forceinline__ VecD cast(const VecS& v) 00090 { 00091 typedef typename VecTraits<VecD>::elem_type D; 00092 return VecTraits<VecD>::make(saturate_cast<D>(v.x), saturate_cast<D>(v.y), saturate_cast<D>(v.z), saturate_cast<D>(v.w)); 00093 } 00094 }; 00095 00096 template <typename VecD, typename VecS> static __device__ __forceinline__ VecD saturate_cast_helper(const VecS& v) 00097 { 00098 return SatCastHelper<VecTraits<VecD>::cn, VecD>::cast(v); 00099 } 00100 } 00101 00102 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar1& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00103 template<typename T> static __device__ __forceinline__ T saturate_cast(const char1& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00104 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort1& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00105 template<typename T> static __device__ __forceinline__ T saturate_cast(const short1& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00106 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint1& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00107 template<typename T> static __device__ __forceinline__ T saturate_cast(const int1& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00108 template<typename T> static __device__ __forceinline__ T saturate_cast(const float1& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00109 template<typename T> static __device__ __forceinline__ T saturate_cast(const double1& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00110 00111 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar2& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00112 template<typename T> static __device__ __forceinline__ T saturate_cast(const char2& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00113 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort2& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00114 template<typename T> static __device__ __forceinline__ T saturate_cast(const short2& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00115 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint2& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00116 template<typename T> static __device__ __forceinline__ T saturate_cast(const int2& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00117 template<typename T> static __device__ __forceinline__ T saturate_cast(const float2& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00118 template<typename T> static __device__ __forceinline__ T saturate_cast(const double2& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00119 00120 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar3& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00121 template<typename T> static __device__ __forceinline__ T saturate_cast(const char3& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00122 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort3& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00123 template<typename T> static __device__ __forceinline__ T saturate_cast(const short3& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00124 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint3& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00125 template<typename T> static __device__ __forceinline__ T saturate_cast(const int3& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00126 template<typename T> static __device__ __forceinline__ T saturate_cast(const float3& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00127 template<typename T> static __device__ __forceinline__ T saturate_cast(const double3& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00128 00129 template<typename T> static __device__ __forceinline__ T saturate_cast(const uchar4& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00130 template<typename T> static __device__ __forceinline__ T saturate_cast(const char4& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00131 template<typename T> static __device__ __forceinline__ T saturate_cast(const ushort4& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00132 template<typename T> static __device__ __forceinline__ T saturate_cast(const short4& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00133 template<typename T> static __device__ __forceinline__ T saturate_cast(const uint4& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00134 template<typename T> static __device__ __forceinline__ T saturate_cast(const int4& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00135 template<typename T> static __device__ __forceinline__ T saturate_cast(const float4& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00136 template<typename T> static __device__ __forceinline__ T saturate_cast(const double4& v) {return vec_math_detail::saturate_cast_helper<T>(v);} 00137 00138 // unary operators 00139 00140 #define CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(op, input_type, output_type) \ 00141 __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a) \ 00142 { \ 00143 return VecTraits<output_type ## 1>::make(op (a.x)); \ 00144 } \ 00145 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a) \ 00146 { \ 00147 return VecTraits<output_type ## 2>::make(op (a.x), op (a.y)); \ 00148 } \ 00149 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a) \ 00150 { \ 00151 return VecTraits<output_type ## 3>::make(op (a.x), op (a.y), op (a.z)); \ 00152 } \ 00153 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a) \ 00154 { \ 00155 return VecTraits<output_type ## 4>::make(op (a.x), op (a.y), op (a.z), op (a.w)); \ 00156 } 00157 00158 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, char, char) 00159 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, short, short) 00160 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, int, int) 00161 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, float, float) 00162 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(-, double, double) 00163 00164 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, uchar, uchar) 00165 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, char, uchar) 00166 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, ushort, uchar) 00167 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, short, uchar) 00168 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, int, uchar) 00169 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, uint, uchar) 00170 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, float, uchar) 00171 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(!, double, uchar) 00172 00173 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, uchar, uchar) 00174 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, char, char) 00175 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, ushort, ushort) 00176 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, short, short) 00177 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, int, int) 00178 CV_CUDEV_IMPLEMENT_VEC_UNARY_OP(~, uint, uint) 00179 00180 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_OP 00181 00182 // unary functions 00183 00184 #define CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(func_name, func, input_type, output_type) \ 00185 __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a) \ 00186 { \ 00187 return VecTraits<output_type ## 1>::make(func (a.x)); \ 00188 } \ 00189 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a) \ 00190 { \ 00191 return VecTraits<output_type ## 2>::make(func (a.x), func (a.y)); \ 00192 } \ 00193 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a) \ 00194 { \ 00195 return VecTraits<output_type ## 3>::make(func (a.x), func (a.y), func (a.z)); \ 00196 } \ 00197 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a) \ 00198 { \ 00199 return VecTraits<output_type ## 4>::make(func (a.x), func (a.y), func (a.z), func (a.w)); \ 00200 } 00201 00202 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, uchar, uchar) 00203 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::abs, char, char) 00204 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, ushort, ushort) 00205 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::abs, short, short) 00206 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::abs, int, int) 00207 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, /*::abs*/, uint, uint) 00208 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::fabsf, float, float) 00209 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(abs, ::fabs, double, double) 00210 00211 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, uchar, float) 00212 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, char, float) 00213 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, ushort, float) 00214 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, short, float) 00215 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, int, float) 00216 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, uint, float) 00217 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrtf, float, float) 00218 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sqrt, ::sqrt, double, double) 00219 00220 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, uchar, float) 00221 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, char, float) 00222 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, ushort, float) 00223 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, short, float) 00224 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, int, float) 00225 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, uint, float) 00226 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::expf, float, float) 00227 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp, ::exp, double, double) 00228 00229 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, uchar, float) 00230 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, char, float) 00231 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, ushort, float) 00232 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, short, float) 00233 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, int, float) 00234 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, uint, float) 00235 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2f, float, float) 00236 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp2, ::exp2, double, double) 00237 00238 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, uchar, float) 00239 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, char, float) 00240 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, ushort, float) 00241 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, short, float) 00242 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, int, float) 00243 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, uint, float) 00244 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10f, float, float) 00245 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(exp10, ::exp10, double, double) 00246 00247 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, uchar, float) 00248 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, char, float) 00249 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, ushort, float) 00250 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, short, float) 00251 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, int, float) 00252 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, uint, float) 00253 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::logf, float, float) 00254 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log, ::log, double, double) 00255 00256 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, uchar, float) 00257 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, char, float) 00258 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, ushort, float) 00259 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, short, float) 00260 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, int, float) 00261 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, uint, float) 00262 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2f, float, float) 00263 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log2, ::log2, double, double) 00264 00265 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, uchar, float) 00266 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, char, float) 00267 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, ushort, float) 00268 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, short, float) 00269 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, int, float) 00270 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, uint, float) 00271 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10f, float, float) 00272 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(log10, ::log10, double, double) 00273 00274 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, uchar, float) 00275 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, char, float) 00276 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, ushort, float) 00277 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, short, float) 00278 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, int, float) 00279 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, uint, float) 00280 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sinf, float, float) 00281 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sin, ::sin, double, double) 00282 00283 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, uchar, float) 00284 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, char, float) 00285 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, ushort, float) 00286 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, short, float) 00287 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, int, float) 00288 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, uint, float) 00289 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cosf, float, float) 00290 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cos, ::cos, double, double) 00291 00292 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, uchar, float) 00293 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, char, float) 00294 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, ushort, float) 00295 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, short, float) 00296 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, int, float) 00297 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, uint, float) 00298 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tanf, float, float) 00299 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tan, ::tan, double, double) 00300 00301 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, uchar, float) 00302 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, char, float) 00303 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, ushort, float) 00304 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, short, float) 00305 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, int, float) 00306 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, uint, float) 00307 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asinf, float, float) 00308 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asin, ::asin, double, double) 00309 00310 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, uchar, float) 00311 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, char, float) 00312 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, ushort, float) 00313 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, short, float) 00314 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, int, float) 00315 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, uint, float) 00316 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acosf, float, float) 00317 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acos, ::acos, double, double) 00318 00319 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, uchar, float) 00320 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, char, float) 00321 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, ushort, float) 00322 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, short, float) 00323 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, int, float) 00324 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, uint, float) 00325 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atanf, float, float) 00326 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atan, ::atan, double, double) 00327 00328 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, uchar, float) 00329 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, char, float) 00330 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, ushort, float) 00331 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, short, float) 00332 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, int, float) 00333 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, uint, float) 00334 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinhf, float, float) 00335 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(sinh, ::sinh, double, double) 00336 00337 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, uchar, float) 00338 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, char, float) 00339 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, ushort, float) 00340 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, short, float) 00341 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, int, float) 00342 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, uint, float) 00343 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::coshf, float, float) 00344 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(cosh, ::cosh, double, double) 00345 00346 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, uchar, float) 00347 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, char, float) 00348 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, ushort, float) 00349 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, short, float) 00350 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, int, float) 00351 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, uint, float) 00352 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanhf, float, float) 00353 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(tanh, ::tanh, double, double) 00354 00355 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, uchar, float) 00356 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, char, float) 00357 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, ushort, float) 00358 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, short, float) 00359 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, int, float) 00360 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, uint, float) 00361 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinhf, float, float) 00362 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(asinh, ::asinh, double, double) 00363 00364 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, uchar, float) 00365 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, char, float) 00366 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, ushort, float) 00367 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, short, float) 00368 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, int, float) 00369 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, uint, float) 00370 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acoshf, float, float) 00371 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(acosh, ::acosh, double, double) 00372 00373 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, uchar, float) 00374 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, char, float) 00375 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, ushort, float) 00376 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, short, float) 00377 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, int, float) 00378 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, uint, float) 00379 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanhf, float, float) 00380 CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC(atanh, ::atanh, double, double) 00381 00382 #undef CV_CUDEV_IMPLEMENT_VEC_UNARY_FUNC 00383 00384 // binary operators (vec & vec) 00385 00386 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(op, input_type, output_type) \ 00387 __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, const input_type ## 1 & b) \ 00388 { \ 00389 return VecTraits<output_type ## 1>::make(a.x op b.x); \ 00390 } \ 00391 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, const input_type ## 2 & b) \ 00392 { \ 00393 return VecTraits<output_type ## 2>::make(a.x op b.x, a.y op b.y); \ 00394 } \ 00395 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, const input_type ## 3 & b) \ 00396 { \ 00397 return VecTraits<output_type ## 3>::make(a.x op b.x, a.y op b.y, a.z op b.z); \ 00398 } \ 00399 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, const input_type ## 4 & b) \ 00400 { \ 00401 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); \ 00402 } 00403 00404 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, uchar, int) 00405 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, char, int) 00406 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, ushort, int) 00407 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, short, int) 00408 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, int, int) 00409 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, uint, uint) 00410 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, float, float) 00411 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(+, double, double) 00412 00413 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uchar, int) 00414 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, char, int) 00415 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, ushort, int) 00416 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, short, int) 00417 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, int, int) 00418 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, uint, uint) 00419 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, float, float) 00420 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(-, double, double) 00421 00422 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, uchar, int) 00423 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, char, int) 00424 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, ushort, int) 00425 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, short, int) 00426 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, int, int) 00427 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, uint, uint) 00428 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, float, float) 00429 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(*, double, double) 00430 00431 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, uchar, int) 00432 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, char, int) 00433 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, ushort, int) 00434 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, short, int) 00435 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, int, int) 00436 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, uint, uint) 00437 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, float, float) 00438 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(/, double, double) 00439 00440 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uchar, uchar) 00441 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, char, uchar) 00442 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, ushort, uchar) 00443 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, short, uchar) 00444 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, int, uchar) 00445 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, uint, uchar) 00446 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, float, uchar) 00447 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(==, double, uchar) 00448 00449 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uchar, uchar) 00450 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, char, uchar) 00451 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, ushort, uchar) 00452 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, short, uchar) 00453 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, int, uchar) 00454 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, uint, uchar) 00455 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, float, uchar) 00456 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(!=, double, uchar) 00457 00458 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, uchar, uchar) 00459 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, char, uchar) 00460 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, ushort, uchar) 00461 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, short, uchar) 00462 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, int, uchar) 00463 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, uint, uchar) 00464 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, float, uchar) 00465 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>, double, uchar) 00466 00467 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, uchar, uchar) 00468 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, char, uchar) 00469 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, ushort, uchar) 00470 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, short, uchar) 00471 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, int, uchar) 00472 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, uint, uchar) 00473 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, float, uchar) 00474 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<, double, uchar) 00475 00476 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, uchar, uchar) 00477 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, char, uchar) 00478 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, ushort, uchar) 00479 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, short, uchar) 00480 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, int, uchar) 00481 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, uint, uchar) 00482 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, float, uchar) 00483 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(>=, double, uchar) 00484 00485 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, uchar, uchar) 00486 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, char, uchar) 00487 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, ushort, uchar) 00488 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, short, uchar) 00489 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, int, uchar) 00490 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, uint, uchar) 00491 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, float, uchar) 00492 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(<=, double, uchar) 00493 00494 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, uchar, uchar) 00495 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, char, uchar) 00496 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, ushort, uchar) 00497 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, short, uchar) 00498 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, int, uchar) 00499 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, uint, uchar) 00500 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, float, uchar) 00501 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&&, double, uchar) 00502 00503 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, uchar, uchar) 00504 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, char, uchar) 00505 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, ushort, uchar) 00506 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, short, uchar) 00507 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, int, uchar) 00508 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, uint, uchar) 00509 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, float, uchar) 00510 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(||, double, uchar) 00511 00512 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, uchar, uchar) 00513 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, char, char) 00514 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, ushort, ushort) 00515 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, short, short) 00516 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, int, int) 00517 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(&, uint, uint) 00518 00519 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, uchar, uchar) 00520 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, char, char) 00521 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, ushort, ushort) 00522 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, short, short) 00523 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, int, int) 00524 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(|, uint, uint) 00525 00526 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uchar, uchar) 00527 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, char, char) 00528 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, ushort, ushort) 00529 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, short, short) 00530 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, int, int) 00531 CV_CUDEV_IMPLEMENT_VEC_BINARY_OP(^, uint, uint) 00532 00533 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_OP 00534 00535 // binary operators (vec & scalar) 00536 00537 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(op, input_type, scalar_type, output_type) \ 00538 __device__ __forceinline__ output_type ## 1 operator op(const input_type ## 1 & a, scalar_type s) \ 00539 { \ 00540 return VecTraits<output_type ## 1>::make(a.x op s); \ 00541 } \ 00542 __device__ __forceinline__ output_type ## 1 operator op(scalar_type s, const input_type ## 1 & b) \ 00543 { \ 00544 return VecTraits<output_type ## 1>::make(s op b.x); \ 00545 } \ 00546 __device__ __forceinline__ output_type ## 2 operator op(const input_type ## 2 & a, scalar_type s) \ 00547 { \ 00548 return VecTraits<output_type ## 2>::make(a.x op s, a.y op s); \ 00549 } \ 00550 __device__ __forceinline__ output_type ## 2 operator op(scalar_type s, const input_type ## 2 & b) \ 00551 { \ 00552 return VecTraits<output_type ## 2>::make(s op b.x, s op b.y); \ 00553 } \ 00554 __device__ __forceinline__ output_type ## 3 operator op(const input_type ## 3 & a, scalar_type s) \ 00555 { \ 00556 return VecTraits<output_type ## 3>::make(a.x op s, a.y op s, a.z op s); \ 00557 } \ 00558 __device__ __forceinline__ output_type ## 3 operator op(scalar_type s, const input_type ## 3 & b) \ 00559 { \ 00560 return VecTraits<output_type ## 3>::make(s op b.x, s op b.y, s op b.z); \ 00561 } \ 00562 __device__ __forceinline__ output_type ## 4 operator op(const input_type ## 4 & a, scalar_type s) \ 00563 { \ 00564 return VecTraits<output_type ## 4>::make(a.x op s, a.y op s, a.z op s, a.w op s); \ 00565 } \ 00566 __device__ __forceinline__ output_type ## 4 operator op(scalar_type s, const input_type ## 4 & b) \ 00567 { \ 00568 return VecTraits<output_type ## 4>::make(s op b.x, s op b.y, s op b.z, s op b.w); \ 00569 } 00570 00571 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, int, int) 00572 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, float, float) 00573 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uchar, double, double) 00574 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, int, int) 00575 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, float, float) 00576 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, char, double, double) 00577 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, int, int) 00578 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, float, float) 00579 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, ushort, double, double) 00580 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, int, int) 00581 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, float, float) 00582 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, short, double, double) 00583 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, int, int) 00584 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, float, float) 00585 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, int, double, double) 00586 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, uint, uint) 00587 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, float, float) 00588 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, uint, double, double) 00589 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, float, float, float) 00590 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, float, double, double) 00591 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(+, double, double, double) 00592 00593 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, int, int) 00594 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, float, float) 00595 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uchar, double, double) 00596 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, int, int) 00597 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, float, float) 00598 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, char, double, double) 00599 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, int, int) 00600 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, float, float) 00601 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, ushort, double, double) 00602 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, int, int) 00603 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, float, float) 00604 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, short, double, double) 00605 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, int, int) 00606 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, float, float) 00607 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, int, double, double) 00608 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, uint, uint) 00609 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, float, float) 00610 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, uint, double, double) 00611 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, float, float, float) 00612 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, float, double, double) 00613 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(-, double, double, double) 00614 00615 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, int, int) 00616 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, float, float) 00617 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uchar, double, double) 00618 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, int, int) 00619 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, float, float) 00620 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, char, double, double) 00621 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, int, int) 00622 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, float, float) 00623 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, ushort, double, double) 00624 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, int, int) 00625 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, float, float) 00626 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, short, double, double) 00627 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, int, int) 00628 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, float, float) 00629 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, int, double, double) 00630 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, uint, uint) 00631 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, float, float) 00632 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, uint, double, double) 00633 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, float, float, float) 00634 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, float, double, double) 00635 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(*, double, double, double) 00636 00637 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, int, int) 00638 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, float, float) 00639 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uchar, double, double) 00640 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, int, int) 00641 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, float, float) 00642 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, char, double, double) 00643 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, int, int) 00644 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, float, float) 00645 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, ushort, double, double) 00646 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, int, int) 00647 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, float, float) 00648 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, short, double, double) 00649 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, int, int) 00650 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, float, float) 00651 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, int, double, double) 00652 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, uint, uint) 00653 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, float, float) 00654 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, uint, double, double) 00655 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, float, float, float) 00656 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, float, double, double) 00657 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(/, double, double, double) 00658 00659 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uchar, uchar, uchar) 00660 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, char, char, uchar) 00661 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, ushort, ushort, uchar) 00662 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, short, short, uchar) 00663 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, int, int, uchar) 00664 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, uint, uint, uchar) 00665 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, float, float, uchar) 00666 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(==, double, double, uchar) 00667 00668 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uchar, uchar, uchar) 00669 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, char, char, uchar) 00670 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, ushort, ushort, uchar) 00671 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, short, short, uchar) 00672 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, int, int, uchar) 00673 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, uint, uint, uchar) 00674 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, float, float, uchar) 00675 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(!=, double, double, uchar) 00676 00677 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, uchar, uchar, uchar) 00678 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, char, char, uchar) 00679 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, ushort, ushort, uchar) 00680 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, short, short, uchar) 00681 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, int, int, uchar) 00682 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, uint, uint, uchar) 00683 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, float, float, uchar) 00684 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>, double, double, uchar) 00685 00686 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, uchar, uchar, uchar) 00687 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, char, char, uchar) 00688 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, ushort, ushort, uchar) 00689 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, short, short, uchar) 00690 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, int, int, uchar) 00691 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, uint, uint, uchar) 00692 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, float, float, uchar) 00693 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<, double, double, uchar) 00694 00695 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, uchar, uchar, uchar) 00696 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, char, char, uchar) 00697 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, ushort, ushort, uchar) 00698 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, short, short, uchar) 00699 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, int, int, uchar) 00700 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, uint, uint, uchar) 00701 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, float, float, uchar) 00702 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(>=, double, double, uchar) 00703 00704 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, uchar, uchar, uchar) 00705 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, char, char, uchar) 00706 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, ushort, ushort, uchar) 00707 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, short, short, uchar) 00708 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, int, int, uchar) 00709 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, uint, uint, uchar) 00710 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, float, float, uchar) 00711 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(<=, double, double, uchar) 00712 00713 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, uchar, uchar, uchar) 00714 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, char, char, uchar) 00715 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, ushort, ushort, uchar) 00716 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, short, short, uchar) 00717 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, int, int, uchar) 00718 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, uint, uint, uchar) 00719 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, float, float, uchar) 00720 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&&, double, double, uchar) 00721 00722 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, uchar, uchar, uchar) 00723 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, char, char, uchar) 00724 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, ushort, ushort, uchar) 00725 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, short, short, uchar) 00726 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, int, int, uchar) 00727 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, uint, uint, uchar) 00728 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, float, float, uchar) 00729 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(||, double, double, uchar) 00730 00731 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, uchar, uchar, uchar) 00732 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, char, char, char) 00733 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, ushort, ushort, ushort) 00734 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, short, short, short) 00735 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, int, int, int) 00736 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(&, uint, uint, uint) 00737 00738 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, uchar, uchar, uchar) 00739 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, char, char, char) 00740 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, ushort, ushort, ushort) 00741 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, short, short, short) 00742 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, int, int, int) 00743 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(|, uint, uint, uint) 00744 00745 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uchar, uchar, uchar) 00746 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, char, char, char) 00747 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, ushort, ushort, ushort) 00748 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, short, short, short) 00749 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, int, int, int) 00750 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP(^, uint, uint, uint) 00751 00752 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_OP 00753 00754 // binary function (vec & vec) 00755 00756 #define CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(func_name, func, input_type, output_type) \ 00757 __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, const input_type ## 1 & b) \ 00758 { \ 00759 return VecTraits<output_type ## 1>::make(func (a.x, b.x)); \ 00760 } \ 00761 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, const input_type ## 2 & b) \ 00762 { \ 00763 return VecTraits<output_type ## 2>::make(func (a.x, b.x), func (a.y, b.y)); \ 00764 } \ 00765 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, const input_type ## 3 & b) \ 00766 { \ 00767 return VecTraits<output_type ## 3>::make(func (a.x, b.x), func (a.y, b.y), func (a.z, b.z)); \ 00768 } \ 00769 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, const input_type ## 4 & b) \ 00770 { \ 00771 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)); \ 00772 } 00773 00774 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, uchar, uchar) 00775 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, char, char) 00776 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, ushort, ushort) 00777 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, short, short) 00778 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, uint, uint) 00779 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::max, int, int) 00780 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::fmaxf, float, float) 00781 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(max, ::fmax, double, double) 00782 00783 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, uchar, uchar) 00784 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, char, char) 00785 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, ushort, ushort) 00786 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, short, short) 00787 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, uint, uint) 00788 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::min, int, int) 00789 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::fminf, float, float) 00790 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(min, ::fmin, double, double) 00791 00792 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, uchar, float) 00793 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, char, float) 00794 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, ushort, float) 00795 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, short, float) 00796 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, uint, float) 00797 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, int, float) 00798 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypotf, float, float) 00799 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(hypot, ::hypot, double, double) 00800 00801 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, uchar, float) 00802 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, char, float) 00803 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, ushort, float) 00804 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, short, float) 00805 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, uint, float) 00806 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, int, float) 00807 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2f, float, float) 00808 CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC(atan2, ::atan2, double, double) 00809 00810 #undef CV_CUDEV_IMPLEMENT_VEC_BINARY_FUNC 00811 00812 // binary function (vec & scalar) 00813 00814 #define CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(func_name, func, input_type, scalar_type, output_type) \ 00815 __device__ __forceinline__ output_type ## 1 func_name(const input_type ## 1 & a, scalar_type s) \ 00816 { \ 00817 return VecTraits<output_type ## 1>::make(func ((output_type) a.x, (output_type) s)); \ 00818 } \ 00819 __device__ __forceinline__ output_type ## 1 func_name(scalar_type s, const input_type ## 1 & b) \ 00820 { \ 00821 return VecTraits<output_type ## 1>::make(func ((output_type) s, (output_type) b.x)); \ 00822 } \ 00823 __device__ __forceinline__ output_type ## 2 func_name(const input_type ## 2 & a, scalar_type s) \ 00824 { \ 00825 return VecTraits<output_type ## 2>::make(func ((output_type) a.x, (output_type) s), func ((output_type) a.y, (output_type) s)); \ 00826 } \ 00827 __device__ __forceinline__ output_type ## 2 func_name(scalar_type s, const input_type ## 2 & b) \ 00828 { \ 00829 return VecTraits<output_type ## 2>::make(func ((output_type) s, (output_type) b.x), func ((output_type) s, (output_type) b.y)); \ 00830 } \ 00831 __device__ __forceinline__ output_type ## 3 func_name(const input_type ## 3 & a, scalar_type s) \ 00832 { \ 00833 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)); \ 00834 } \ 00835 __device__ __forceinline__ output_type ## 3 func_name(scalar_type s, const input_type ## 3 & b) \ 00836 { \ 00837 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)); \ 00838 } \ 00839 __device__ __forceinline__ output_type ## 4 func_name(const input_type ## 4 & a, scalar_type s) \ 00840 { \ 00841 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)); \ 00842 } \ 00843 __device__ __forceinline__ output_type ## 4 func_name(scalar_type s, const input_type ## 4 & b) \ 00844 { \ 00845 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)); \ 00846 } 00847 00848 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, uchar, uchar, uchar) 00849 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, uchar, float, float) 00850 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, uchar, double, double) 00851 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, char, char, char) 00852 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, char, float, float) 00853 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, char, double, double) 00854 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, ushort, ushort, ushort) 00855 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, ushort, float, float) 00856 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, ushort, double, double) 00857 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, short, short, short) 00858 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, short, float, float) 00859 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, short, double, double) 00860 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, uint, uint, uint) 00861 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, uint, float, float) 00862 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, uint, double, double) 00863 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::max, int, int, int) 00864 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, int, float, float) 00865 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, int, double, double) 00866 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmaxf, float, float, float) 00867 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, float, double, double) 00868 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(max, ::fmax, double, double, double) 00869 00870 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, uchar, uchar, uchar) 00871 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, uchar, float, float) 00872 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, uchar, double, double) 00873 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, char, char, char) 00874 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, char, float, float) 00875 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, char, double, double) 00876 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, ushort, ushort, ushort) 00877 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, ushort, float, float) 00878 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, ushort, double, double) 00879 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, short, short, short) 00880 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, short, float, float) 00881 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, short, double, double) 00882 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, uint, uint, uint) 00883 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, uint, float, float) 00884 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, uint, double, double) 00885 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::min, int, int, int) 00886 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, int, float, float) 00887 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, int, double, double) 00888 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fminf, float, float, float) 00889 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, float, double, double) 00890 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(min, ::fmin, double, double, double) 00891 00892 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, uchar, float, float) 00893 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, uchar, double, double) 00894 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, char, float, float) 00895 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, char, double, double) 00896 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, ushort, float, float) 00897 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, ushort, double, double) 00898 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, short, float, float) 00899 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, short, double, double) 00900 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, uint, float, float) 00901 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, uint, double, double) 00902 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, int, float, float) 00903 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, int, double, double) 00904 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypotf, float, float, float) 00905 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, float, double, double) 00906 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(hypot, ::hypot, double, double, double) 00907 00908 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, uchar, float, float) 00909 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, uchar, double, double) 00910 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, char, float, float) 00911 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, char, double, double) 00912 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, ushort, float, float) 00913 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, ushort, double, double) 00914 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, short, float, float) 00915 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, short, double, double) 00916 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, uint, float, float) 00917 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, uint, double, double) 00918 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, int, float, float) 00919 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, int, double, double) 00920 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2f, float, float, float) 00921 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, float, double, double) 00922 CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC(atan2, ::atan2, double, double, double) 00923 00924 #undef CV_CUDEV_IMPLEMENT_SCALAR_BINARY_FUNC 00925 00926 }}} // namespace cv { namespace cuda { namespace device 00927 00928 //! @endcond 00929 00930 #endif // __OPENCV_CUDA_VECMATH_HPP__ 00931
Generated on Tue Jul 12 2022 14:47:51 by
1.7.2