1 /*M/////////////////////////////////////////////////////////////////////////////////////// 2 // 3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. 4 // 5 // By downloading, copying, installing or using the software you agree to this license. 6 // If you do not agree to this license, do not download, install, 7 // copy or use the software. 8 // 9 // 10 // License Agreement 11 // For Open Source Computer Vision Library 12 // 13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved. 14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved. 15 // Third party copyrights are property of their respective owners. 16 // 17 // Redistribution and use in source and binary forms, with or without modification, 18 // are permitted provided that the following conditions are met: 19 // 20 // * Redistribution's of source code must retain the above copyright notice, 21 // this list of conditions and the following disclaimer. 22 // 23 // * Redistribution's in binary form must reproduce the above copyright notice, 24 // this list of conditions and the following disclaimer in the documentation 25 // and/or other materials provided with the distribution. 26 // 27 // * The name of the copyright holders may not be used to endorse or promote products 28 // derived from this software without specific prior written permission. 29 // 30 // This software is provided by the copyright holders and contributors "as is" and 31 // any express or implied warranties, including, but not limited to, the implied 32 // warranties of merchantability and fitness for a particular purpose are disclaimed. 33 // In no event shall the Intel Corporation or contributors be liable for any direct, 34 // indirect, incidental, special, exemplary, or consequential damages 35 // (including, but not limited to, procurement of substitute goods or services; 36 // loss of use, data, or profits; or business interruption) however caused 37 // and on any theory of liability, whether in contract, strict liability, 38 // or tort (including negligence or otherwise) arising in any way out of 39 // the use of this software, even if advised of the possibility of such damage. 40 // 41 //M*/ 42 43 #ifndef OPENCV_CUDA_FUNCTIONAL_HPP 44 #define OPENCV_CUDA_FUNCTIONAL_HPP 45 46 #include <functional> 47 #include "saturate_cast.hpp" 48 #include "vec_traits.hpp" 49 #include "type_traits.hpp" 50 51 /** @file 52 * @deprecated Use @ref cudev instead. 53 */ 54 55 //! @cond IGNORED 56 57 namespace cv { namespace cuda { namespace device 58 { 59 // Function Objects 60 template<typename Argument, typename Result> struct unary_function 61 { 62 typedef Argument argument_type; 63 typedef Result result_type; 64 }; 65 template<typename Argument1, typename Argument2, typename Result> struct binary_function 66 { 67 typedef Argument1 first_argument_type; 68 typedef Argument2 second_argument_type; 69 typedef Result result_type; 70 }; 71 72 // Arithmetic Operations 73 template <typename T> struct plus : binary_function<T, T, T> 74 { operator ()cv::cuda::device::plus75 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 76 typename TypeTraits<T>::ParameterType b) const 77 { 78 return a + b; 79 } pluscv::cuda::device::plus80 __host__ __device__ __forceinline__ plus() {} pluscv::cuda::device::plus81 __host__ __device__ __forceinline__ plus(const plus&) {} 82 }; 83 84 template <typename T> struct minus : binary_function<T, T, T> 85 { operator ()cv::cuda::device::minus86 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 87 typename TypeTraits<T>::ParameterType b) const 88 { 89 return a - b; 90 } minuscv::cuda::device::minus91 __host__ __device__ __forceinline__ minus() {} minuscv::cuda::device::minus92 __host__ __device__ __forceinline__ minus(const minus&) {} 93 }; 94 95 template <typename T> struct multiplies : binary_function<T, T, T> 96 { operator ()cv::cuda::device::multiplies97 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 98 typename TypeTraits<T>::ParameterType b) const 99 { 100 return a * b; 101 } multipliescv::cuda::device::multiplies102 __host__ __device__ __forceinline__ multiplies() {} multipliescv::cuda::device::multiplies103 __host__ __device__ __forceinline__ multiplies(const multiplies&) {} 104 }; 105 106 template <typename T> struct divides : binary_function<T, T, T> 107 { operator ()cv::cuda::device::divides108 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 109 typename TypeTraits<T>::ParameterType b) const 110 { 111 return a / b; 112 } dividescv::cuda::device::divides113 __host__ __device__ __forceinline__ divides() {} dividescv::cuda::device::divides114 __host__ __device__ __forceinline__ divides(const divides&) {} 115 }; 116 117 template <typename T> struct modulus : binary_function<T, T, T> 118 { operator ()cv::cuda::device::modulus119 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 120 typename TypeTraits<T>::ParameterType b) const 121 { 122 return a % b; 123 } moduluscv::cuda::device::modulus124 __host__ __device__ __forceinline__ modulus() {} moduluscv::cuda::device::modulus125 __host__ __device__ __forceinline__ modulus(const modulus&) {} 126 }; 127 128 template <typename T> struct negate : unary_function<T, T> 129 { operator ()cv::cuda::device::negate130 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a) const 131 { 132 return -a; 133 } negatecv::cuda::device::negate134 __host__ __device__ __forceinline__ negate() {} negatecv::cuda::device::negate135 __host__ __device__ __forceinline__ negate(const negate&) {} 136 }; 137 138 // Comparison Operations 139 template <typename T> struct equal_to : binary_function<T, T, bool> 140 { operator ()cv::cuda::device::equal_to141 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 142 typename TypeTraits<T>::ParameterType b) const 143 { 144 return a == b; 145 } equal_tocv::cuda::device::equal_to146 __host__ __device__ __forceinline__ equal_to() {} equal_tocv::cuda::device::equal_to147 __host__ __device__ __forceinline__ equal_to(const equal_to&) {} 148 }; 149 150 template <typename T> struct not_equal_to : binary_function<T, T, bool> 151 { operator ()cv::cuda::device::not_equal_to152 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 153 typename TypeTraits<T>::ParameterType b) const 154 { 155 return a != b; 156 } not_equal_tocv::cuda::device::not_equal_to157 __host__ __device__ __forceinline__ not_equal_to() {} not_equal_tocv::cuda::device::not_equal_to158 __host__ __device__ __forceinline__ not_equal_to(const not_equal_to&) {} 159 }; 160 161 template <typename T> struct greater : binary_function<T, T, bool> 162 { operator ()cv::cuda::device::greater163 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 164 typename TypeTraits<T>::ParameterType b) const 165 { 166 return a > b; 167 } greatercv::cuda::device::greater168 __host__ __device__ __forceinline__ greater() {} greatercv::cuda::device::greater169 __host__ __device__ __forceinline__ greater(const greater&) {} 170 }; 171 172 template <typename T> struct less : binary_function<T, T, bool> 173 { operator ()cv::cuda::device::less174 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 175 typename TypeTraits<T>::ParameterType b) const 176 { 177 return a < b; 178 } lesscv::cuda::device::less179 __host__ __device__ __forceinline__ less() {} lesscv::cuda::device::less180 __host__ __device__ __forceinline__ less(const less&) {} 181 }; 182 183 template <typename T> struct greater_equal : binary_function<T, T, bool> 184 { operator ()cv::cuda::device::greater_equal185 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 186 typename TypeTraits<T>::ParameterType b) const 187 { 188 return a >= b; 189 } greater_equalcv::cuda::device::greater_equal190 __host__ __device__ __forceinline__ greater_equal() {} greater_equalcv::cuda::device::greater_equal191 __host__ __device__ __forceinline__ greater_equal(const greater_equal&) {} 192 }; 193 194 template <typename T> struct less_equal : binary_function<T, T, bool> 195 { operator ()cv::cuda::device::less_equal196 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 197 typename TypeTraits<T>::ParameterType b) const 198 { 199 return a <= b; 200 } less_equalcv::cuda::device::less_equal201 __host__ __device__ __forceinline__ less_equal() {} less_equalcv::cuda::device::less_equal202 __host__ __device__ __forceinline__ less_equal(const less_equal&) {} 203 }; 204 205 // Logical Operations 206 template <typename T> struct logical_and : binary_function<T, T, bool> 207 { operator ()cv::cuda::device::logical_and208 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 209 typename TypeTraits<T>::ParameterType b) const 210 { 211 return a && b; 212 } logical_andcv::cuda::device::logical_and213 __host__ __device__ __forceinline__ logical_and() {} logical_andcv::cuda::device::logical_and214 __host__ __device__ __forceinline__ logical_and(const logical_and&) {} 215 }; 216 217 template <typename T> struct logical_or : binary_function<T, T, bool> 218 { operator ()cv::cuda::device::logical_or219 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 220 typename TypeTraits<T>::ParameterType b) const 221 { 222 return a || b; 223 } logical_orcv::cuda::device::logical_or224 __host__ __device__ __forceinline__ logical_or() {} logical_orcv::cuda::device::logical_or225 __host__ __device__ __forceinline__ logical_or(const logical_or&) {} 226 }; 227 228 template <typename T> struct logical_not : unary_function<T, bool> 229 { operator ()cv::cuda::device::logical_not230 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a) const 231 { 232 return !a; 233 } logical_notcv::cuda::device::logical_not234 __host__ __device__ __forceinline__ logical_not() {} logical_notcv::cuda::device::logical_not235 __host__ __device__ __forceinline__ logical_not(const logical_not&) {} 236 }; 237 238 // Bitwise Operations 239 template <typename T> struct bit_and : binary_function<T, T, T> 240 { operator ()cv::cuda::device::bit_and241 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 242 typename TypeTraits<T>::ParameterType b) const 243 { 244 return a & b; 245 } bit_andcv::cuda::device::bit_and246 __host__ __device__ __forceinline__ bit_and() {} bit_andcv::cuda::device::bit_and247 __host__ __device__ __forceinline__ bit_and(const bit_and&) {} 248 }; 249 250 template <typename T> struct bit_or : binary_function<T, T, T> 251 { operator ()cv::cuda::device::bit_or252 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 253 typename TypeTraits<T>::ParameterType b) const 254 { 255 return a | b; 256 } bit_orcv::cuda::device::bit_or257 __host__ __device__ __forceinline__ bit_or() {} bit_orcv::cuda::device::bit_or258 __host__ __device__ __forceinline__ bit_or(const bit_or&) {} 259 }; 260 261 template <typename T> struct bit_xor : binary_function<T, T, T> 262 { operator ()cv::cuda::device::bit_xor263 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 264 typename TypeTraits<T>::ParameterType b) const 265 { 266 return a ^ b; 267 } bit_xorcv::cuda::device::bit_xor268 __host__ __device__ __forceinline__ bit_xor() {} bit_xorcv::cuda::device::bit_xor269 __host__ __device__ __forceinline__ bit_xor(const bit_xor&) {} 270 }; 271 272 template <typename T> struct bit_not : unary_function<T, T> 273 { operator ()cv::cuda::device::bit_not274 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType v) const 275 { 276 return ~v; 277 } bit_notcv::cuda::device::bit_not278 __host__ __device__ __forceinline__ bit_not() {} bit_notcv::cuda::device::bit_not279 __host__ __device__ __forceinline__ bit_not(const bit_not&) {} 280 }; 281 282 // Generalized Identity Operations 283 template <typename T> struct identity : unary_function<T, T> 284 { operator ()cv::cuda::device::identity285 __device__ __forceinline__ typename TypeTraits<T>::ParameterType operator()(typename TypeTraits<T>::ParameterType x) const 286 { 287 return x; 288 } identitycv::cuda::device::identity289 __host__ __device__ __forceinline__ identity() {} identitycv::cuda::device::identity290 __host__ __device__ __forceinline__ identity(const identity&) {} 291 }; 292 293 template <typename T1, typename T2> struct project1st : binary_function<T1, T2, T1> 294 { operator ()cv::cuda::device::project1st295 __device__ __forceinline__ typename TypeTraits<T1>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const 296 { 297 return lhs; 298 } project1stcv::cuda::device::project1st299 __host__ __device__ __forceinline__ project1st() {} project1stcv::cuda::device::project1st300 __host__ __device__ __forceinline__ project1st(const project1st&) {} 301 }; 302 303 template <typename T1, typename T2> struct project2nd : binary_function<T1, T2, T2> 304 { operator ()cv::cuda::device::project2nd305 __device__ __forceinline__ typename TypeTraits<T2>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const 306 { 307 return rhs; 308 } project2ndcv::cuda::device::project2nd309 __host__ __device__ __forceinline__ project2nd() {} project2ndcv::cuda::device::project2nd310 __host__ __device__ __forceinline__ project2nd(const project2nd&) {} 311 }; 312 313 // Min/Max Operations 314 315 #define OPENCV_CUDA_IMPLEMENT_MINMAX(name, type, op) \ 316 template <> struct name<type> : binary_function<type, type, type> \ 317 { \ 318 __device__ __forceinline__ type operator()(type lhs, type rhs) const {return op(lhs, rhs);} \ 319 __host__ __device__ __forceinline__ name() {}\ 320 __host__ __device__ __forceinline__ name(const name&) {}\ 321 }; 322 323 template <typename T> struct maximum : binary_function<T, T, T> 324 { operator ()cv::cuda::device::maximum325 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType lhs, typename TypeTraits<T>::ParameterType rhs) const 326 { 327 return max(lhs, rhs); 328 } maximumcv::cuda::device::maximum329 __host__ __device__ __forceinline__ maximum() {} maximumcv::cuda::device::maximum330 __host__ __device__ __forceinline__ maximum(const maximum&) {} 331 }; 332 333 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, uchar, ::max) 334 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, schar, ::max) 335 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, char, ::max) 336 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, ushort, ::max) 337 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, short, ::max) 338 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, int, ::max) 339 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, uint, ::max) 340 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, float, ::fmax) 341 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, double, ::fmax) 342 343 template <typename T> struct minimum : binary_function<T, T, T> 344 { operator ()cv::cuda::device::minimum345 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType lhs, typename TypeTraits<T>::ParameterType rhs) const 346 { 347 return min(lhs, rhs); 348 } minimumcv::cuda::device::minimum349 __host__ __device__ __forceinline__ minimum() {} minimumcv::cuda::device::minimum350 __host__ __device__ __forceinline__ minimum(const minimum&) {} 351 }; 352 353 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, uchar, ::min) 354 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, schar, ::min) 355 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, char, ::min) 356 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, ushort, ::min) 357 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, short, ::min) 358 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, int, ::min) 359 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, uint, ::min) 360 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, float, ::fmin) 361 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, double, ::fmin) 362 363 #undef OPENCV_CUDA_IMPLEMENT_MINMAX 364 365 // Math functions 366 367 template <typename T> struct abs_func : unary_function<T, T> 368 { operator ()cv::cuda::device::abs_func369 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType x) const 370 { 371 return abs(x); 372 } 373 abs_funccv::cuda::device::abs_func374 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func375 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 376 }; 377 template <> struct abs_func<unsigned char> : unary_function<unsigned char, unsigned char> 378 { operator ()cv::cuda::device::abs_func379 __device__ __forceinline__ unsigned char operator ()(unsigned char x) const 380 { 381 return x; 382 } 383 abs_funccv::cuda::device::abs_func384 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func385 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 386 }; 387 template <> struct abs_func<signed char> : unary_function<signed char, signed char> 388 { operator ()cv::cuda::device::abs_func389 __device__ __forceinline__ signed char operator ()(signed char x) const 390 { 391 return ::abs((int)x); 392 } 393 abs_funccv::cuda::device::abs_func394 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func395 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 396 }; 397 template <> struct abs_func<char> : unary_function<char, char> 398 { operator ()cv::cuda::device::abs_func399 __device__ __forceinline__ char operator ()(char x) const 400 { 401 return ::abs((int)x); 402 } 403 abs_funccv::cuda::device::abs_func404 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func405 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 406 }; 407 template <> struct abs_func<unsigned short> : unary_function<unsigned short, unsigned short> 408 { operator ()cv::cuda::device::abs_func409 __device__ __forceinline__ unsigned short operator ()(unsigned short x) const 410 { 411 return x; 412 } 413 abs_funccv::cuda::device::abs_func414 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func415 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 416 }; 417 template <> struct abs_func<short> : unary_function<short, short> 418 { operator ()cv::cuda::device::abs_func419 __device__ __forceinline__ short operator ()(short x) const 420 { 421 return ::abs((int)x); 422 } 423 abs_funccv::cuda::device::abs_func424 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func425 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 426 }; 427 template <> struct abs_func<unsigned int> : unary_function<unsigned int, unsigned int> 428 { operator ()cv::cuda::device::abs_func429 __device__ __forceinline__ unsigned int operator ()(unsigned int x) const 430 { 431 return x; 432 } 433 abs_funccv::cuda::device::abs_func434 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func435 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 436 }; 437 template <> struct abs_func<int> : unary_function<int, int> 438 { operator ()cv::cuda::device::abs_func439 __device__ __forceinline__ int operator ()(int x) const 440 { 441 return ::abs(x); 442 } 443 abs_funccv::cuda::device::abs_func444 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func445 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 446 }; 447 template <> struct abs_func<float> : unary_function<float, float> 448 { operator ()cv::cuda::device::abs_func449 __device__ __forceinline__ float operator ()(float x) const 450 { 451 return ::fabsf(x); 452 } 453 abs_funccv::cuda::device::abs_func454 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func455 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 456 }; 457 template <> struct abs_func<double> : unary_function<double, double> 458 { operator ()cv::cuda::device::abs_func459 __device__ __forceinline__ double operator ()(double x) const 460 { 461 return ::fabs(x); 462 } 463 abs_funccv::cuda::device::abs_func464 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func465 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 466 }; 467 468 #define OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(name, func) \ 469 template <typename T> struct name ## _func : unary_function<T, float> \ 470 { \ 471 __device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v) const \ 472 { \ 473 return func ## f(v); \ 474 } \ 475 __host__ __device__ __forceinline__ name ## _func() {} \ 476 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 477 }; \ 478 template <> struct name ## _func<double> : unary_function<double, double> \ 479 { \ 480 __device__ __forceinline__ double operator ()(double v) const \ 481 { \ 482 return func(v); \ 483 } \ 484 __host__ __device__ __forceinline__ name ## _func() {} \ 485 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 486 }; 487 488 #define OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(name, func) \ 489 template <typename T> struct name ## _func : binary_function<T, T, float> \ 490 { \ 491 __device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v1, typename TypeTraits<T>::ParameterType v2) const \ 492 { \ 493 return func ## f(v1, v2); \ 494 } \ 495 __host__ __device__ __forceinline__ name ## _func() {} \ 496 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 497 }; \ 498 template <> struct name ## _func<double> : binary_function<double, double, double> \ 499 { \ 500 __device__ __forceinline__ double operator ()(double v1, double v2) const \ 501 { \ 502 return func(v1, v2); \ 503 } \ 504 __host__ __device__ __forceinline__ name ## _func() {} \ 505 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 506 }; 507 508 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sqrt, ::sqrt) 509 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp, ::exp) 510 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp2, ::exp2) 511 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp10, ::exp10) 512 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log, ::log) 513 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log2, ::log2) 514 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log10, ::log10) 515 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sin, ::sin) 516 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(cos, ::cos) 517 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(tan, ::tan) 518 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(asin, ::asin) 519 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(acos, ::acos) 520 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(atan, ::atan) 521 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sinh, ::sinh) 522 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(cosh, ::cosh) 523 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(tanh, ::tanh) 524 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(asinh, ::asinh) 525 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(acosh, ::acosh) 526 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(atanh, ::atanh) 527 528 OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(hypot, ::hypot) 529 OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(atan2, ::atan2) 530 OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(pow, ::pow) 531 532 #undef OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR 533 #undef OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR_NO_DOUBLE 534 #undef OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR 535 536 template<typename T> struct hypot_sqr_func : binary_function<T, T, float> 537 { operator ()cv::cuda::device::hypot_sqr_func538 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType src1, typename TypeTraits<T>::ParameterType src2) const 539 { 540 return src1 * src1 + src2 * src2; 541 } hypot_sqr_funccv::cuda::device::hypot_sqr_func542 __host__ __device__ __forceinline__ hypot_sqr_func() {} hypot_sqr_funccv::cuda::device::hypot_sqr_func543 __host__ __device__ __forceinline__ hypot_sqr_func(const hypot_sqr_func&) {} 544 }; 545 546 // Saturate Cast Functor 547 template <typename T, typename D> struct saturate_cast_func : unary_function<T, D> 548 { operator ()cv::cuda::device::saturate_cast_func549 __device__ __forceinline__ D operator ()(typename TypeTraits<T>::ParameterType v) const 550 { 551 return saturate_cast<D>(v); 552 } saturate_cast_funccv::cuda::device::saturate_cast_func553 __host__ __device__ __forceinline__ saturate_cast_func() {} saturate_cast_funccv::cuda::device::saturate_cast_func554 __host__ __device__ __forceinline__ saturate_cast_func(const saturate_cast_func&) {} 555 }; 556 557 // Threshold Functors 558 template <typename T> struct thresh_binary_func : unary_function<T, T> 559 { thresh_binary_funccv::cuda::device::thresh_binary_func560 __host__ __device__ __forceinline__ thresh_binary_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {} 561 operator ()cv::cuda::device::thresh_binary_func562 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 563 { 564 return (src > thresh) * maxVal; 565 } 566 thresh_binary_funccv::cuda::device::thresh_binary_func567 __host__ __device__ __forceinline__ thresh_binary_func() {} thresh_binary_funccv::cuda::device::thresh_binary_func568 __host__ __device__ __forceinline__ thresh_binary_func(const thresh_binary_func& other) 569 : thresh(other.thresh), maxVal(other.maxVal) {} 570 571 T thresh; 572 T maxVal; 573 }; 574 575 template <typename T> struct thresh_binary_inv_func : unary_function<T, T> 576 { thresh_binary_inv_funccv::cuda::device::thresh_binary_inv_func577 __host__ __device__ __forceinline__ thresh_binary_inv_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {} 578 operator ()cv::cuda::device::thresh_binary_inv_func579 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 580 { 581 return (src <= thresh) * maxVal; 582 } 583 thresh_binary_inv_funccv::cuda::device::thresh_binary_inv_func584 __host__ __device__ __forceinline__ thresh_binary_inv_func() {} thresh_binary_inv_funccv::cuda::device::thresh_binary_inv_func585 __host__ __device__ __forceinline__ thresh_binary_inv_func(const thresh_binary_inv_func& other) 586 : thresh(other.thresh), maxVal(other.maxVal) {} 587 588 T thresh; 589 T maxVal; 590 }; 591 592 template <typename T> struct thresh_trunc_func : unary_function<T, T> 593 { thresh_trunc_funccv::cuda::device::thresh_trunc_func594 explicit __host__ __device__ __forceinline__ thresh_trunc_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {CV_UNUSED(maxVal_);} 595 operator ()cv::cuda::device::thresh_trunc_func596 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 597 { 598 return minimum<T>()(src, thresh); 599 } 600 thresh_trunc_funccv::cuda::device::thresh_trunc_func601 __host__ __device__ __forceinline__ thresh_trunc_func() {} thresh_trunc_funccv::cuda::device::thresh_trunc_func602 __host__ __device__ __forceinline__ thresh_trunc_func(const thresh_trunc_func& other) 603 : thresh(other.thresh) {} 604 605 T thresh; 606 }; 607 608 template <typename T> struct thresh_to_zero_func : unary_function<T, T> 609 { thresh_to_zero_funccv::cuda::device::thresh_to_zero_func610 explicit __host__ __device__ __forceinline__ thresh_to_zero_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {CV_UNUSED(maxVal_);} 611 operator ()cv::cuda::device::thresh_to_zero_func612 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 613 { 614 return (src > thresh) * src; 615 } 616 thresh_to_zero_funccv::cuda::device::thresh_to_zero_func617 __host__ __device__ __forceinline__ thresh_to_zero_func() {} thresh_to_zero_funccv::cuda::device::thresh_to_zero_func618 __host__ __device__ __forceinline__ thresh_to_zero_func(const thresh_to_zero_func& other) 619 : thresh(other.thresh) {} 620 621 T thresh; 622 }; 623 624 template <typename T> struct thresh_to_zero_inv_func : unary_function<T, T> 625 { thresh_to_zero_inv_funccv::cuda::device::thresh_to_zero_inv_func626 explicit __host__ __device__ __forceinline__ thresh_to_zero_inv_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {CV_UNUSED(maxVal_);} 627 operator ()cv::cuda::device::thresh_to_zero_inv_func628 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 629 { 630 return (src <= thresh) * src; 631 } 632 thresh_to_zero_inv_funccv::cuda::device::thresh_to_zero_inv_func633 __host__ __device__ __forceinline__ thresh_to_zero_inv_func() {} thresh_to_zero_inv_funccv::cuda::device::thresh_to_zero_inv_func634 __host__ __device__ __forceinline__ thresh_to_zero_inv_func(const thresh_to_zero_inv_func& other) 635 : thresh(other.thresh) {} 636 637 T thresh; 638 }; 639 640 // Function Object Adaptors 641 template <typename Predicate> struct unary_negate : unary_function<typename Predicate::argument_type, bool> 642 { unary_negatecv::cuda::device::unary_negate643 explicit __host__ __device__ __forceinline__ unary_negate(const Predicate& p) : pred(p) {} 644 operator ()cv::cuda::device::unary_negate645 __device__ __forceinline__ bool operator()(typename TypeTraits<typename Predicate::argument_type>::ParameterType x) const 646 { 647 return !pred(x); 648 } 649 unary_negatecv::cuda::device::unary_negate650 __host__ __device__ __forceinline__ unary_negate() {} unary_negatecv::cuda::device::unary_negate651 __host__ __device__ __forceinline__ unary_negate(const unary_negate& other) : pred(other.pred) {} 652 653 Predicate pred; 654 }; 655 not1(const Predicate & pred)656 template <typename Predicate> __host__ __device__ __forceinline__ unary_negate<Predicate> not1(const Predicate& pred) 657 { 658 return unary_negate<Predicate>(pred); 659 } 660 661 template <typename Predicate> struct binary_negate : binary_function<typename Predicate::first_argument_type, typename Predicate::second_argument_type, bool> 662 { binary_negatecv::cuda::device::binary_negate663 explicit __host__ __device__ __forceinline__ binary_negate(const Predicate& p) : pred(p) {} 664 operator ()cv::cuda::device::binary_negate665 __device__ __forceinline__ bool operator()(typename TypeTraits<typename Predicate::first_argument_type>::ParameterType x, 666 typename TypeTraits<typename Predicate::second_argument_type>::ParameterType y) const 667 { 668 return !pred(x,y); 669 } 670 binary_negatecv::cuda::device::binary_negate671 __host__ __device__ __forceinline__ binary_negate() {} binary_negatecv::cuda::device::binary_negate672 __host__ __device__ __forceinline__ binary_negate(const binary_negate& other) : pred(other.pred) {} 673 674 Predicate pred; 675 }; 676 not2(const BinaryPredicate & pred)677 template <typename BinaryPredicate> __host__ __device__ __forceinline__ binary_negate<BinaryPredicate> not2(const BinaryPredicate& pred) 678 { 679 return binary_negate<BinaryPredicate>(pred); 680 } 681 682 template <typename Op> struct binder1st : unary_function<typename Op::second_argument_type, typename Op::result_type> 683 { binder1stcv::cuda::device::binder1st684 __host__ __device__ __forceinline__ binder1st(const Op& op_, const typename Op::first_argument_type& arg1_) : op(op_), arg1(arg1_) {} 685 operator ()cv::cuda::device::binder1st686 __device__ __forceinline__ typename Op::result_type operator ()(typename TypeTraits<typename Op::second_argument_type>::ParameterType a) const 687 { 688 return op(arg1, a); 689 } 690 binder1stcv::cuda::device::binder1st691 __host__ __device__ __forceinline__ binder1st() {} binder1stcv::cuda::device::binder1st692 __host__ __device__ __forceinline__ binder1st(const binder1st& other) : op(other.op), arg1(other.arg1) {} 693 694 Op op; 695 typename Op::first_argument_type arg1; 696 }; 697 bind1st(const Op & op,const T & x)698 template <typename Op, typename T> __host__ __device__ __forceinline__ binder1st<Op> bind1st(const Op& op, const T& x) 699 { 700 return binder1st<Op>(op, typename Op::first_argument_type(x)); 701 } 702 703 template <typename Op> struct binder2nd : unary_function<typename Op::first_argument_type, typename Op::result_type> 704 { binder2ndcv::cuda::device::binder2nd705 __host__ __device__ __forceinline__ binder2nd(const Op& op_, const typename Op::second_argument_type& arg2_) : op(op_), arg2(arg2_) {} 706 operator ()cv::cuda::device::binder2nd707 __forceinline__ __device__ typename Op::result_type operator ()(typename TypeTraits<typename Op::first_argument_type>::ParameterType a) const 708 { 709 return op(a, arg2); 710 } 711 binder2ndcv::cuda::device::binder2nd712 __host__ __device__ __forceinline__ binder2nd() {} binder2ndcv::cuda::device::binder2nd713 __host__ __device__ __forceinline__ binder2nd(const binder2nd& other) : op(other.op), arg2(other.arg2) {} 714 715 Op op; 716 typename Op::second_argument_type arg2; 717 }; 718 bind2nd(const Op & op,const T & x)719 template <typename Op, typename T> __host__ __device__ __forceinline__ binder2nd<Op> bind2nd(const Op& op, const T& x) 720 { 721 return binder2nd<Op>(op, typename Op::second_argument_type(x)); 722 } 723 724 // Functor Traits 725 template <typename F> struct IsUnaryFunction 726 { 727 typedef char Yes; 728 struct No {Yes a[2];}; 729 730 template <typename T, typename D> static Yes check(unary_function<T, D>); 731 static No check(...); 732 733 static F makeF(); 734 735 enum { value = (sizeof(check(makeF())) == sizeof(Yes)) }; 736 }; 737 738 template <typename F> struct IsBinaryFunction 739 { 740 typedef char Yes; 741 struct No {Yes a[2];}; 742 743 template <typename T1, typename T2, typename D> static Yes check(binary_function<T1, T2, D>); 744 static No check(...); 745 746 static F makeF(); 747 748 enum { value = (sizeof(check(makeF())) == sizeof(Yes)) }; 749 }; 750 751 namespace functional_detail 752 { 753 template <size_t src_elem_size, size_t dst_elem_size> struct UnOpShift { enum { shift = 1 }; }; 754 template <size_t src_elem_size> struct UnOpShift<src_elem_size, 1> { enum { shift = 4 }; }; 755 template <size_t src_elem_size> struct UnOpShift<src_elem_size, 2> { enum { shift = 2 }; }; 756 757 template <typename T, typename D> struct DefaultUnaryShift 758 { 759 enum { shift = UnOpShift<sizeof(T), sizeof(D)>::shift }; 760 }; 761 762 template <size_t src_elem_size1, size_t src_elem_size2, size_t dst_elem_size> struct BinOpShift { enum { shift = 1 }; }; 763 template <size_t src_elem_size1, size_t src_elem_size2> struct BinOpShift<src_elem_size1, src_elem_size2, 1> { enum { shift = 4 }; }; 764 template <size_t src_elem_size1, size_t src_elem_size2> struct BinOpShift<src_elem_size1, src_elem_size2, 2> { enum { shift = 2 }; }; 765 766 template <typename T1, typename T2, typename D> struct DefaultBinaryShift 767 { 768 enum { shift = BinOpShift<sizeof(T1), sizeof(T2), sizeof(D)>::shift }; 769 }; 770 771 template <typename Func, bool unary = IsUnaryFunction<Func>::value> struct ShiftDispatcher; 772 template <typename Func> struct ShiftDispatcher<Func, true> 773 { 774 enum { shift = DefaultUnaryShift<typename Func::argument_type, typename Func::result_type>::shift }; 775 }; 776 template <typename Func> struct ShiftDispatcher<Func, false> 777 { 778 enum { shift = DefaultBinaryShift<typename Func::first_argument_type, typename Func::second_argument_type, typename Func::result_type>::shift }; 779 }; 780 } 781 782 template <typename Func> struct DefaultTransformShift 783 { 784 enum { shift = functional_detail::ShiftDispatcher<Func>::shift }; 785 }; 786 787 template <typename Func> struct DefaultTransformFunctorTraits 788 { 789 enum { simple_block_dim_x = 16 }; 790 enum { simple_block_dim_y = 16 }; 791 792 enum { smart_block_dim_x = 16 }; 793 enum { smart_block_dim_y = 16 }; 794 enum { smart_shift = DefaultTransformShift<Func>::shift }; 795 }; 796 797 template <typename Func> struct TransformFunctorTraits : DefaultTransformFunctorTraits<Func> {}; 798 799 #define OPENCV_CUDA_TRANSFORM_FUNCTOR_TRAITS(type) \ 800 template <> struct TransformFunctorTraits< type > : DefaultTransformFunctorTraits< type > 801 }}} // namespace cv { namespace cuda { namespace cudev 802 803 //! @endcond 804 805 #endif // OPENCV_CUDA_FUNCTIONAL_HPP 806