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
functional.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_FUNCTIONAL_HPP__ 00044 #define __OPENCV_CUDA_FUNCTIONAL_HPP__ 00045 00046 #include <functional> 00047 #include "saturate_cast.hpp " 00048 #include "vec_traits.hpp " 00049 #include "type_traits.hpp " 00050 #include "device_functions.h" 00051 00052 /** @file 00053 * @deprecated Use @ref cudev instead. 00054 */ 00055 00056 //! @cond IGNORED 00057 00058 namespace cv { namespace cuda { namespace device 00059 { 00060 // Function Objects 00061 template<typename Argument, typename Result> struct unary_function : public std::unary_function<Argument, Result> {}; 00062 template<typename Argument1, typename Argument2, typename Result> struct binary_function : public std::binary_function<Argument1, Argument2, Result> {}; 00063 00064 // Arithmetic Operations 00065 template <typename T> struct plus : binary_function<T, T, T> 00066 { 00067 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 00068 typename TypeTraits<T>::ParameterType b) const 00069 { 00070 return a + b; 00071 } 00072 __host__ __device__ __forceinline__ plus() {} 00073 __host__ __device__ __forceinline__ plus(const plus&) {} 00074 }; 00075 00076 template <typename T> struct minus : binary_function<T, T, T> 00077 { 00078 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 00079 typename TypeTraits<T>::ParameterType b) const 00080 { 00081 return a - b; 00082 } 00083 __host__ __device__ __forceinline__ minus() {} 00084 __host__ __device__ __forceinline__ minus(const minus&) {} 00085 }; 00086 00087 template <typename T> struct multiplies : binary_function<T, T, T> 00088 { 00089 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 00090 typename TypeTraits<T>::ParameterType b) const 00091 { 00092 return a * b; 00093 } 00094 __host__ __device__ __forceinline__ multiplies() {} 00095 __host__ __device__ __forceinline__ multiplies(const multiplies&) {} 00096 }; 00097 00098 template <typename T> struct divides : binary_function<T, T, T> 00099 { 00100 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 00101 typename TypeTraits<T>::ParameterType b) const 00102 { 00103 return a / b; 00104 } 00105 __host__ __device__ __forceinline__ divides() {} 00106 __host__ __device__ __forceinline__ divides(const divides&) {} 00107 }; 00108 00109 template <typename T> struct modulus : binary_function<T, T, T> 00110 { 00111 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 00112 typename TypeTraits<T>::ParameterType b) const 00113 { 00114 return a % b; 00115 } 00116 __host__ __device__ __forceinline__ modulus() {} 00117 __host__ __device__ __forceinline__ modulus(const modulus&) {} 00118 }; 00119 00120 template <typename T> struct negate : unary_function<T, T> 00121 { 00122 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a) const 00123 { 00124 return -a; 00125 } 00126 __host__ __device__ __forceinline__ negate() {} 00127 __host__ __device__ __forceinline__ negate(const negate&) {} 00128 }; 00129 00130 // Comparison Operations 00131 template <typename T> struct equal_to : binary_function<T, T, bool> 00132 { 00133 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 00134 typename TypeTraits<T>::ParameterType b) const 00135 { 00136 return a == b; 00137 } 00138 __host__ __device__ __forceinline__ equal_to() {} 00139 __host__ __device__ __forceinline__ equal_to(const equal_to&) {} 00140 }; 00141 00142 template <typename T> struct not_equal_to : binary_function<T, T, bool> 00143 { 00144 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 00145 typename TypeTraits<T>::ParameterType b) const 00146 { 00147 return a != b; 00148 } 00149 __host__ __device__ __forceinline__ not_equal_to() {} 00150 __host__ __device__ __forceinline__ not_equal_to(const not_equal_to&) {} 00151 }; 00152 00153 template <typename T> struct greater : binary_function<T, T, bool> 00154 { 00155 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 00156 typename TypeTraits<T>::ParameterType b) const 00157 { 00158 return a > b; 00159 } 00160 __host__ __device__ __forceinline__ greater() {} 00161 __host__ __device__ __forceinline__ greater(const greater&) {} 00162 }; 00163 00164 template <typename T> struct less : binary_function<T, T, bool> 00165 { 00166 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 00167 typename TypeTraits<T>::ParameterType b) const 00168 { 00169 return a < b; 00170 } 00171 __host__ __device__ __forceinline__ less() {} 00172 __host__ __device__ __forceinline__ less(const less&) {} 00173 }; 00174 00175 template <typename T> struct greater_equal : binary_function<T, T, bool> 00176 { 00177 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 00178 typename TypeTraits<T>::ParameterType b) const 00179 { 00180 return a >= b; 00181 } 00182 __host__ __device__ __forceinline__ greater_equal() {} 00183 __host__ __device__ __forceinline__ greater_equal(const greater_equal&) {} 00184 }; 00185 00186 template <typename T> struct less_equal : binary_function<T, T, bool> 00187 { 00188 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 00189 typename TypeTraits<T>::ParameterType b) const 00190 { 00191 return a <= b; 00192 } 00193 __host__ __device__ __forceinline__ less_equal() {} 00194 __host__ __device__ __forceinline__ less_equal(const less_equal&) {} 00195 }; 00196 00197 // Logical Operations 00198 template <typename T> struct logical_and : binary_function<T, T, bool> 00199 { 00200 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 00201 typename TypeTraits<T>::ParameterType b) const 00202 { 00203 return a && b; 00204 } 00205 __host__ __device__ __forceinline__ logical_and() {} 00206 __host__ __device__ __forceinline__ logical_and(const logical_and&) {} 00207 }; 00208 00209 template <typename T> struct logical_or : binary_function<T, T, bool> 00210 { 00211 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 00212 typename TypeTraits<T>::ParameterType b) const 00213 { 00214 return a || b; 00215 } 00216 __host__ __device__ __forceinline__ logical_or() {} 00217 __host__ __device__ __forceinline__ logical_or(const logical_or&) {} 00218 }; 00219 00220 template <typename T> struct logical_not : unary_function<T, bool> 00221 { 00222 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a) const 00223 { 00224 return !a; 00225 } 00226 __host__ __device__ __forceinline__ logical_not() {} 00227 __host__ __device__ __forceinline__ logical_not(const logical_not&) {} 00228 }; 00229 00230 // Bitwise Operations 00231 template <typename T> struct bit_and : binary_function<T, T, T> 00232 { 00233 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 00234 typename TypeTraits<T>::ParameterType b) const 00235 { 00236 return a & b; 00237 } 00238 __host__ __device__ __forceinline__ bit_and() {} 00239 __host__ __device__ __forceinline__ bit_and(const bit_and&) {} 00240 }; 00241 00242 template <typename T> struct bit_or : binary_function<T, T, T> 00243 { 00244 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 00245 typename TypeTraits<T>::ParameterType b) const 00246 { 00247 return a | b; 00248 } 00249 __host__ __device__ __forceinline__ bit_or() {} 00250 __host__ __device__ __forceinline__ bit_or(const bit_or&) {} 00251 }; 00252 00253 template <typename T> struct bit_xor : binary_function<T, T, T> 00254 { 00255 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 00256 typename TypeTraits<T>::ParameterType b) const 00257 { 00258 return a ^ b; 00259 } 00260 __host__ __device__ __forceinline__ bit_xor() {} 00261 __host__ __device__ __forceinline__ bit_xor(const bit_xor&) {} 00262 }; 00263 00264 template <typename T> struct bit_not : unary_function<T, T> 00265 { 00266 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType v) const 00267 { 00268 return ~v; 00269 } 00270 __host__ __device__ __forceinline__ bit_not() {} 00271 __host__ __device__ __forceinline__ bit_not(const bit_not&) {} 00272 }; 00273 00274 // Generalized Identity Operations 00275 template <typename T> struct identity : unary_function<T, T> 00276 { 00277 __device__ __forceinline__ typename TypeTraits<T>::ParameterType operator()(typename TypeTraits<T>::ParameterType x) const 00278 { 00279 return x; 00280 } 00281 __host__ __device__ __forceinline__ identity() {} 00282 __host__ __device__ __forceinline__ identity(const identity&) {} 00283 }; 00284 00285 template <typename T1, typename T2> struct project1st : binary_function<T1, T2, T1> 00286 { 00287 __device__ __forceinline__ typename TypeTraits<T1>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const 00288 { 00289 return lhs; 00290 } 00291 __host__ __device__ __forceinline__ project1st() {} 00292 __host__ __device__ __forceinline__ project1st(const project1st&) {} 00293 }; 00294 00295 template <typename T1, typename T2> struct project2nd : binary_function<T1, T2, T2> 00296 { 00297 __device__ __forceinline__ typename TypeTraits<T2>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const 00298 { 00299 return rhs; 00300 } 00301 __host__ __device__ __forceinline__ project2nd() {} 00302 __host__ __device__ __forceinline__ project2nd(const project2nd&) {} 00303 }; 00304 00305 // Min/Max Operations 00306 00307 #define OPENCV_CUDA_IMPLEMENT_MINMAX(name, type, op) \ 00308 template <> struct name<type> : binary_function<type, type, type> \ 00309 { \ 00310 __device__ __forceinline__ type operator()(type lhs, type rhs) const {return op(lhs, rhs);} \ 00311 __host__ __device__ __forceinline__ name() {}\ 00312 __host__ __device__ __forceinline__ name(const name&) {}\ 00313 }; 00314 00315 template <typename T> struct maximum : binary_function<T, T, T> 00316 { 00317 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType lhs, typename TypeTraits<T>::ParameterType rhs) const 00318 { 00319 return max(lhs, rhs); 00320 } 00321 __host__ __device__ __forceinline__ maximum() {} 00322 __host__ __device__ __forceinline__ maximum(const maximum&) {} 00323 }; 00324 00325 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, uchar, ::max) 00326 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, schar, ::max) 00327 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, char, ::max) 00328 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, ushort, ::max) 00329 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, short, ::max) 00330 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, int, ::max) 00331 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, uint, ::max) 00332 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, float, ::fmax) 00333 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, double, ::fmax) 00334 00335 template <typename T> struct minimum : binary_function<T, T, T> 00336 { 00337 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType lhs, typename TypeTraits<T>::ParameterType rhs) const 00338 { 00339 return min(lhs, rhs); 00340 } 00341 __host__ __device__ __forceinline__ minimum() {} 00342 __host__ __device__ __forceinline__ minimum(const minimum&) {} 00343 }; 00344 00345 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, uchar, ::min) 00346 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, schar, ::min) 00347 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, char, ::min) 00348 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, ushort, ::min) 00349 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, short, ::min) 00350 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, int, ::min) 00351 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, uint, ::min) 00352 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, float, ::fmin) 00353 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, double, ::fmin) 00354 00355 #undef OPENCV_CUDA_IMPLEMENT_MINMAX 00356 00357 // Math functions 00358 00359 template <typename T> struct abs_func : unary_function<T, T> 00360 { 00361 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType x) const 00362 { 00363 return abs(x); 00364 } 00365 00366 __host__ __device__ __forceinline__ abs_func() {} 00367 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 00368 }; 00369 template <> struct abs_func<unsigned char> : unary_function<unsigned char, unsigned char> 00370 { 00371 __device__ __forceinline__ unsigned char operator ()(unsigned char x) const 00372 { 00373 return x; 00374 } 00375 00376 __host__ __device__ __forceinline__ abs_func() {} 00377 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 00378 }; 00379 template <> struct abs_func<signed char> : unary_function<signed char, signed char> 00380 { 00381 __device__ __forceinline__ signed char operator ()(signed char x) const 00382 { 00383 return ::abs((int)x); 00384 } 00385 00386 __host__ __device__ __forceinline__ abs_func() {} 00387 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 00388 }; 00389 template <> struct abs_func<char> : unary_function<char, char> 00390 { 00391 __device__ __forceinline__ char operator ()(char x) const 00392 { 00393 return ::abs((int)x); 00394 } 00395 00396 __host__ __device__ __forceinline__ abs_func() {} 00397 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 00398 }; 00399 template <> struct abs_func<unsigned short> : unary_function<unsigned short, unsigned short> 00400 { 00401 __device__ __forceinline__ unsigned short operator ()(unsigned short x) const 00402 { 00403 return x; 00404 } 00405 00406 __host__ __device__ __forceinline__ abs_func() {} 00407 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 00408 }; 00409 template <> struct abs_func<short> : unary_function<short, short> 00410 { 00411 __device__ __forceinline__ short operator ()(short x) const 00412 { 00413 return ::abs((int)x); 00414 } 00415 00416 __host__ __device__ __forceinline__ abs_func() {} 00417 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 00418 }; 00419 template <> struct abs_func<unsigned int> : unary_function<unsigned int, unsigned int> 00420 { 00421 __device__ __forceinline__ unsigned int operator ()(unsigned int x) const 00422 { 00423 return x; 00424 } 00425 00426 __host__ __device__ __forceinline__ abs_func() {} 00427 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 00428 }; 00429 template <> struct abs_func<int> : unary_function<int, int> 00430 { 00431 __device__ __forceinline__ int operator ()(int x) const 00432 { 00433 return ::abs(x); 00434 } 00435 00436 __host__ __device__ __forceinline__ abs_func() {} 00437 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 00438 }; 00439 template <> struct abs_func<float> : unary_function<float, float> 00440 { 00441 __device__ __forceinline__ float operator ()(float x) const 00442 { 00443 return ::fabsf(x); 00444 } 00445 00446 __host__ __device__ __forceinline__ abs_func() {} 00447 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 00448 }; 00449 template <> struct abs_func<double> : unary_function<double, double> 00450 { 00451 __device__ __forceinline__ double operator ()(double x) const 00452 { 00453 return ::fabs(x); 00454 } 00455 00456 __host__ __device__ __forceinline__ abs_func() {} 00457 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 00458 }; 00459 00460 #define OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(name, func) \ 00461 template <typename T> struct name ## _func : unary_function<T, float> \ 00462 { \ 00463 __device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v) const \ 00464 { \ 00465 return func ## f(v); \ 00466 } \ 00467 __host__ __device__ __forceinline__ name ## _func() {} \ 00468 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 00469 }; \ 00470 template <> struct name ## _func<double> : unary_function<double, double> \ 00471 { \ 00472 __device__ __forceinline__ double operator ()(double v) const \ 00473 { \ 00474 return func(v); \ 00475 } \ 00476 __host__ __device__ __forceinline__ name ## _func() {} \ 00477 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 00478 }; 00479 00480 #define OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(name, func) \ 00481 template <typename T> struct name ## _func : binary_function<T, T, float> \ 00482 { \ 00483 __device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v1, typename TypeTraits<T>::ParameterType v2) const \ 00484 { \ 00485 return func ## f(v1, v2); \ 00486 } \ 00487 __host__ __device__ __forceinline__ name ## _func() {} \ 00488 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 00489 }; \ 00490 template <> struct name ## _func<double> : binary_function<double, double, double> \ 00491 { \ 00492 __device__ __forceinline__ double operator ()(double v1, double v2) const \ 00493 { \ 00494 return func(v1, v2); \ 00495 } \ 00496 __host__ __device__ __forceinline__ name ## _func() {} \ 00497 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 00498 }; 00499 00500 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sqrt, ::sqrt) 00501 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp, ::exp) 00502 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp2, ::exp2) 00503 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp10, ::exp10) 00504 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log, ::log) 00505 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log2, ::log2) 00506 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log10, ::log10) 00507 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sin, ::sin) 00508 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(cos, ::cos) 00509 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(tan, ::tan) 00510 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(asin, ::asin) 00511 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(acos, ::acos) 00512 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(atan, ::atan) 00513 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sinh, ::sinh) 00514 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(cosh, ::cosh) 00515 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(tanh, ::tanh) 00516 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(asinh, ::asinh) 00517 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(acosh, ::acosh) 00518 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(atanh, ::atanh) 00519 00520 OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(hypot, ::hypot) 00521 OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(atan2, ::atan2) 00522 OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(pow, ::pow) 00523 00524 #undef OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR 00525 #undef OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR_NO_DOUBLE 00526 #undef OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR 00527 00528 template<typename T> struct hypot_sqr_func : binary_function<T, T, float> 00529 { 00530 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType src1, typename TypeTraits<T>::ParameterType src2) const 00531 { 00532 return src1 * src1 + src2 * src2; 00533 } 00534 __host__ __device__ __forceinline__ hypot_sqr_func() {} 00535 __host__ __device__ __forceinline__ hypot_sqr_func(const hypot_sqr_func&) {} 00536 }; 00537 00538 // Saturate Cast Functor 00539 template <typename T, typename D> struct saturate_cast_func : unary_function<T, D> 00540 { 00541 __device__ __forceinline__ D operator ()(typename TypeTraits<T>::ParameterType v) const 00542 { 00543 return saturate_cast<D>(v); 00544 } 00545 __host__ __device__ __forceinline__ saturate_cast_func() {} 00546 __host__ __device__ __forceinline__ saturate_cast_func(const saturate_cast_func&) {} 00547 }; 00548 00549 // Threshold Functors 00550 template <typename T> struct thresh_binary_func : unary_function<T, T> 00551 { 00552 __host__ __device__ __forceinline__ thresh_binary_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {} 00553 00554 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 00555 { 00556 return (src > thresh) * maxVal; 00557 } 00558 00559 __host__ __device__ __forceinline__ thresh_binary_func() {} 00560 __host__ __device__ __forceinline__ thresh_binary_func(const thresh_binary_func& other) 00561 : thresh(other.thresh), maxVal(other.maxVal) {} 00562 00563 T thresh; 00564 T maxVal; 00565 }; 00566 00567 template <typename T> struct thresh_binary_inv_func : unary_function<T, T> 00568 { 00569 __host__ __device__ __forceinline__ thresh_binary_inv_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {} 00570 00571 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 00572 { 00573 return (src <= thresh) * maxVal; 00574 } 00575 00576 __host__ __device__ __forceinline__ thresh_binary_inv_func() {} 00577 __host__ __device__ __forceinline__ thresh_binary_inv_func(const thresh_binary_inv_func& other) 00578 : thresh(other.thresh), maxVal(other.maxVal) {} 00579 00580 T thresh; 00581 T maxVal; 00582 }; 00583 00584 template <typename T> struct thresh_trunc_func : unary_function<T, T> 00585 { 00586 explicit __host__ __device__ __forceinline__ thresh_trunc_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;} 00587 00588 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 00589 { 00590 return minimum<T>()(src, thresh); 00591 } 00592 00593 __host__ __device__ __forceinline__ thresh_trunc_func() {} 00594 __host__ __device__ __forceinline__ thresh_trunc_func(const thresh_trunc_func& other) 00595 : thresh(other.thresh) {} 00596 00597 T thresh; 00598 }; 00599 00600 template <typename T> struct thresh_to_zero_func : unary_function<T, T> 00601 { 00602 explicit __host__ __device__ __forceinline__ thresh_to_zero_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;} 00603 00604 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 00605 { 00606 return (src > thresh) * src; 00607 } 00608 00609 __host__ __device__ __forceinline__ thresh_to_zero_func() {} 00610 __host__ __device__ __forceinline__ thresh_to_zero_func(const thresh_to_zero_func& other) 00611 : thresh(other.thresh) {} 00612 00613 T thresh; 00614 }; 00615 00616 template <typename T> struct thresh_to_zero_inv_func : unary_function<T, T> 00617 { 00618 explicit __host__ __device__ __forceinline__ thresh_to_zero_inv_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;} 00619 00620 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 00621 { 00622 return (src <= thresh) * src; 00623 } 00624 00625 __host__ __device__ __forceinline__ thresh_to_zero_inv_func() {} 00626 __host__ __device__ __forceinline__ thresh_to_zero_inv_func(const thresh_to_zero_inv_func& other) 00627 : thresh(other.thresh) {} 00628 00629 T thresh; 00630 }; 00631 00632 // Function Object Adaptors 00633 template <typename Predicate> struct unary_negate : unary_function<typename Predicate::argument_type, bool> 00634 { 00635 explicit __host__ __device__ __forceinline__ unary_negate(const Predicate& p) : pred(p) {} 00636 00637 __device__ __forceinline__ bool operator()(typename TypeTraits<typename Predicate::argument_type>::ParameterType x) const 00638 { 00639 return !pred(x); 00640 } 00641 00642 __host__ __device__ __forceinline__ unary_negate() {} 00643 __host__ __device__ __forceinline__ unary_negate(const unary_negate& other) : pred(other.pred) {} 00644 00645 Predicate pred; 00646 }; 00647 00648 template <typename Predicate> __host__ __device__ __forceinline__ unary_negate<Predicate> not1(const Predicate& pred) 00649 { 00650 return unary_negate<Predicate>(pred); 00651 } 00652 00653 template <typename Predicate> struct binary_negate : binary_function<typename Predicate::first_argument_type, typename Predicate::second_argument_type, bool> 00654 { 00655 explicit __host__ __device__ __forceinline__ binary_negate(const Predicate& p) : pred(p) {} 00656 00657 __device__ __forceinline__ bool operator()(typename TypeTraits<typename Predicate::first_argument_type>::ParameterType x, 00658 typename TypeTraits<typename Predicate::second_argument_type>::ParameterType y) const 00659 { 00660 return !pred(x,y); 00661 } 00662 00663 __host__ __device__ __forceinline__ binary_negate() {} 00664 __host__ __device__ __forceinline__ binary_negate(const binary_negate& other) : pred(other.pred) {} 00665 00666 Predicate pred; 00667 }; 00668 00669 template <typename BinaryPredicate> __host__ __device__ __forceinline__ binary_negate<BinaryPredicate> not2(const BinaryPredicate& pred) 00670 { 00671 return binary_negate<BinaryPredicate>(pred); 00672 } 00673 00674 template <typename Op> struct binder1st : unary_function<typename Op::second_argument_type, typename Op::result_type> 00675 { 00676 __host__ __device__ __forceinline__ binder1st(const Op& op_, const typename Op::first_argument_type& arg1_) : op(op_), arg1(arg1_) {} 00677 00678 __device__ __forceinline__ typename Op::result_type operator ()(typename TypeTraits<typename Op::second_argument_type>::ParameterType a) const 00679 { 00680 return op(arg1, a); 00681 } 00682 00683 __host__ __device__ __forceinline__ binder1st() {} 00684 __host__ __device__ __forceinline__ binder1st(const binder1st& other) : op(other.op), arg1(other.arg1) {} 00685 00686 Op op; 00687 typename Op::first_argument_type arg1; 00688 }; 00689 00690 template <typename Op, typename T> __host__ __device__ __forceinline__ binder1st<Op> bind1st(const Op& op, const T& x) 00691 { 00692 return binder1st<Op>(op, typename Op::first_argument_type(x)); 00693 } 00694 00695 template <typename Op> struct binder2nd : unary_function<typename Op::first_argument_type, typename Op::result_type> 00696 { 00697 __host__ __device__ __forceinline__ binder2nd(const Op& op_, const typename Op::second_argument_type& arg2_) : op(op_), arg2(arg2_) {} 00698 00699 __forceinline__ __device__ typename Op::result_type operator ()(typename TypeTraits<typename Op::first_argument_type>::ParameterType a) const 00700 { 00701 return op(a, arg2); 00702 } 00703 00704 __host__ __device__ __forceinline__ binder2nd() {} 00705 __host__ __device__ __forceinline__ binder2nd(const binder2nd& other) : op(other.op), arg2(other.arg2) {} 00706 00707 Op op; 00708 typename Op::second_argument_type arg2; 00709 }; 00710 00711 template <typename Op, typename T> __host__ __device__ __forceinline__ binder2nd<Op> bind2nd(const Op& op, const T& x) 00712 { 00713 return binder2nd<Op>(op, typename Op::second_argument_type(x)); 00714 } 00715 00716 // Functor Traits 00717 template <typename F> struct IsUnaryFunction 00718 { 00719 typedef char Yes; 00720 struct No {Yes a[2];}; 00721 00722 template <typename T, typename D> static Yes check(unary_function<T, D>); 00723 static No check(...); 00724 00725 static F makeF(); 00726 00727 enum { value = (sizeof(check(makeF())) == sizeof(Yes)) }; 00728 }; 00729 00730 template <typename F> struct IsBinaryFunction 00731 { 00732 typedef char Yes; 00733 struct No {Yes a[2];}; 00734 00735 template <typename T1, typename T2, typename D> static Yes check(binary_function<T1, T2, D>); 00736 static No check(...); 00737 00738 static F makeF(); 00739 00740 enum { value = (sizeof(check(makeF())) == sizeof(Yes)) }; 00741 }; 00742 00743 namespace functional_detail 00744 { 00745 template <size_t src_elem_size, size_t dst_elem_size> struct UnOpShift { enum { shift = 1 }; }; 00746 template <size_t src_elem_size> struct UnOpShift<src_elem_size, 1> { enum { shift = 4 }; }; 00747 template <size_t src_elem_size> struct UnOpShift<src_elem_size, 2> { enum { shift = 2 }; }; 00748 00749 template <typename T, typename D> struct DefaultUnaryShift 00750 { 00751 enum { shift = UnOpShift<sizeof(T), sizeof(D)>::shift }; 00752 }; 00753 00754 template <size_t src_elem_size1, size_t src_elem_size2, size_t dst_elem_size> struct BinOpShift { enum { shift = 1 }; }; 00755 template <size_t src_elem_size1, size_t src_elem_size2> struct BinOpShift<src_elem_size1, src_elem_size2, 1> { enum { shift = 4 }; }; 00756 template <size_t src_elem_size1, size_t src_elem_size2> struct BinOpShift<src_elem_size1, src_elem_size2, 2> { enum { shift = 2 }; }; 00757 00758 template <typename T1, typename T2, typename D> struct DefaultBinaryShift 00759 { 00760 enum { shift = BinOpShift<sizeof(T1), sizeof(T2), sizeof(D)>::shift }; 00761 }; 00762 00763 template <typename Func, bool unary = IsUnaryFunction<Func>::value> struct ShiftDispatcher; 00764 template <typename Func> struct ShiftDispatcher<Func, true> 00765 { 00766 enum { shift = DefaultUnaryShift<typename Func::argument_type, typename Func::result_type>::shift }; 00767 }; 00768 template <typename Func> struct ShiftDispatcher<Func, false> 00769 { 00770 enum { shift = DefaultBinaryShift<typename Func::first_argument_type, typename Func::second_argument_type, typename Func::result_type>::shift }; 00771 }; 00772 } 00773 00774 template <typename Func> struct DefaultTransformShift 00775 { 00776 enum { shift = functional_detail::ShiftDispatcher<Func>::shift }; 00777 }; 00778 00779 template <typename Func> struct DefaultTransformFunctorTraits 00780 { 00781 enum { simple_block_dim_x = 16 }; 00782 enum { simple_block_dim_y = 16 }; 00783 00784 enum { smart_block_dim_x = 16 }; 00785 enum { smart_block_dim_y = 16 }; 00786 enum { smart_shift = DefaultTransformShift<Func>::shift }; 00787 }; 00788 00789 template <typename Func> struct TransformFunctorTraits : DefaultTransformFunctorTraits<Func> {}; 00790 00791 #define OPENCV_CUDA_TRANSFORM_FUNCTOR_TRAITS(type) \ 00792 template <> struct TransformFunctorTraits< type > : DefaultTransformFunctorTraits< type > 00793 }}} // namespace cv { namespace cuda { namespace cudev 00794 00795 //! @endcond 00796 00797 #endif // __OPENCV_CUDA_FUNCTIONAL_HPP__ 00798
Generated on Tue Jul 12 2022 14:46:43 by
1.7.2