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 #include "device_functions.h" 51 52 /** @file 53 * @deprecated Use @ref cudev instead. 54 */ 55 56 //! @cond IGNORED 57 58 namespace cv { namespace cuda { namespace device 59 { 60 // Function Objects 61 template<typename Argument, typename Result> struct unary_function : public std::unary_function<Argument, Result> {}; 62 template<typename Argument1, typename Argument2, typename Result> struct binary_function : public std::binary_function<Argument1, Argument2, Result> {}; 63 64 // Arithmetic Operations 65 template <typename T> struct plus : binary_function<T, T, T> 66 { operator ()cv::cuda::device::plus67 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 68 typename TypeTraits<T>::ParameterType b) const 69 { 70 return a + b; 71 } pluscv::cuda::device::plus72 __host__ __device__ __forceinline__ plus() {} pluscv::cuda::device::plus73 __host__ __device__ __forceinline__ plus(const plus&) {} 74 }; 75 76 template <typename T> struct minus : binary_function<T, T, T> 77 { operator ()cv::cuda::device::minus78 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 79 typename TypeTraits<T>::ParameterType b) const 80 { 81 return a - b; 82 } minuscv::cuda::device::minus83 __host__ __device__ __forceinline__ minus() {} minuscv::cuda::device::minus84 __host__ __device__ __forceinline__ minus(const minus&) {} 85 }; 86 87 template <typename T> struct multiplies : binary_function<T, T, T> 88 { operator ()cv::cuda::device::multiplies89 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 90 typename TypeTraits<T>::ParameterType b) const 91 { 92 return a * b; 93 } multipliescv::cuda::device::multiplies94 __host__ __device__ __forceinline__ multiplies() {} multipliescv::cuda::device::multiplies95 __host__ __device__ __forceinline__ multiplies(const multiplies&) {} 96 }; 97 98 template <typename T> struct divides : binary_function<T, T, T> 99 { operator ()cv::cuda::device::divides100 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 101 typename TypeTraits<T>::ParameterType b) const 102 { 103 return a / b; 104 } dividescv::cuda::device::divides105 __host__ __device__ __forceinline__ divides() {} dividescv::cuda::device::divides106 __host__ __device__ __forceinline__ divides(const divides&) {} 107 }; 108 109 template <typename T> struct modulus : binary_function<T, T, T> 110 { operator ()cv::cuda::device::modulus111 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 112 typename TypeTraits<T>::ParameterType b) const 113 { 114 return a % b; 115 } moduluscv::cuda::device::modulus116 __host__ __device__ __forceinline__ modulus() {} moduluscv::cuda::device::modulus117 __host__ __device__ __forceinline__ modulus(const modulus&) {} 118 }; 119 120 template <typename T> struct negate : unary_function<T, T> 121 { operator ()cv::cuda::device::negate122 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a) const 123 { 124 return -a; 125 } negatecv::cuda::device::negate126 __host__ __device__ __forceinline__ negate() {} negatecv::cuda::device::negate127 __host__ __device__ __forceinline__ negate(const negate&) {} 128 }; 129 130 // Comparison Operations 131 template <typename T> struct equal_to : binary_function<T, T, bool> 132 { operator ()cv::cuda::device::equal_to133 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 134 typename TypeTraits<T>::ParameterType b) const 135 { 136 return a == b; 137 } equal_tocv::cuda::device::equal_to138 __host__ __device__ __forceinline__ equal_to() {} equal_tocv::cuda::device::equal_to139 __host__ __device__ __forceinline__ equal_to(const equal_to&) {} 140 }; 141 142 template <typename T> struct not_equal_to : binary_function<T, T, bool> 143 { operator ()cv::cuda::device::not_equal_to144 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 145 typename TypeTraits<T>::ParameterType b) const 146 { 147 return a != b; 148 } not_equal_tocv::cuda::device::not_equal_to149 __host__ __device__ __forceinline__ not_equal_to() {} not_equal_tocv::cuda::device::not_equal_to150 __host__ __device__ __forceinline__ not_equal_to(const not_equal_to&) {} 151 }; 152 153 template <typename T> struct greater : binary_function<T, T, bool> 154 { operator ()cv::cuda::device::greater155 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 156 typename TypeTraits<T>::ParameterType b) const 157 { 158 return a > b; 159 } greatercv::cuda::device::greater160 __host__ __device__ __forceinline__ greater() {} greatercv::cuda::device::greater161 __host__ __device__ __forceinline__ greater(const greater&) {} 162 }; 163 164 template <typename T> struct less : binary_function<T, T, bool> 165 { operator ()cv::cuda::device::less166 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 167 typename TypeTraits<T>::ParameterType b) const 168 { 169 return a < b; 170 } lesscv::cuda::device::less171 __host__ __device__ __forceinline__ less() {} lesscv::cuda::device::less172 __host__ __device__ __forceinline__ less(const less&) {} 173 }; 174 175 template <typename T> struct greater_equal : binary_function<T, T, bool> 176 { operator ()cv::cuda::device::greater_equal177 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 178 typename TypeTraits<T>::ParameterType b) const 179 { 180 return a >= b; 181 } greater_equalcv::cuda::device::greater_equal182 __host__ __device__ __forceinline__ greater_equal() {} greater_equalcv::cuda::device::greater_equal183 __host__ __device__ __forceinline__ greater_equal(const greater_equal&) {} 184 }; 185 186 template <typename T> struct less_equal : binary_function<T, T, bool> 187 { operator ()cv::cuda::device::less_equal188 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 189 typename TypeTraits<T>::ParameterType b) const 190 { 191 return a <= b; 192 } less_equalcv::cuda::device::less_equal193 __host__ __device__ __forceinline__ less_equal() {} less_equalcv::cuda::device::less_equal194 __host__ __device__ __forceinline__ less_equal(const less_equal&) {} 195 }; 196 197 // Logical Operations 198 template <typename T> struct logical_and : binary_function<T, T, bool> 199 { operator ()cv::cuda::device::logical_and200 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 201 typename TypeTraits<T>::ParameterType b) const 202 { 203 return a && b; 204 } logical_andcv::cuda::device::logical_and205 __host__ __device__ __forceinline__ logical_and() {} logical_andcv::cuda::device::logical_and206 __host__ __device__ __forceinline__ logical_and(const logical_and&) {} 207 }; 208 209 template <typename T> struct logical_or : binary_function<T, T, bool> 210 { operator ()cv::cuda::device::logical_or211 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a, 212 typename TypeTraits<T>::ParameterType b) const 213 { 214 return a || b; 215 } logical_orcv::cuda::device::logical_or216 __host__ __device__ __forceinline__ logical_or() {} logical_orcv::cuda::device::logical_or217 __host__ __device__ __forceinline__ logical_or(const logical_or&) {} 218 }; 219 220 template <typename T> struct logical_not : unary_function<T, bool> 221 { operator ()cv::cuda::device::logical_not222 __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a) const 223 { 224 return !a; 225 } logical_notcv::cuda::device::logical_not226 __host__ __device__ __forceinline__ logical_not() {} logical_notcv::cuda::device::logical_not227 __host__ __device__ __forceinline__ logical_not(const logical_not&) {} 228 }; 229 230 // Bitwise Operations 231 template <typename T> struct bit_and : binary_function<T, T, T> 232 { operator ()cv::cuda::device::bit_and233 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 234 typename TypeTraits<T>::ParameterType b) const 235 { 236 return a & b; 237 } bit_andcv::cuda::device::bit_and238 __host__ __device__ __forceinline__ bit_and() {} bit_andcv::cuda::device::bit_and239 __host__ __device__ __forceinline__ bit_and(const bit_and&) {} 240 }; 241 242 template <typename T> struct bit_or : binary_function<T, T, T> 243 { operator ()cv::cuda::device::bit_or244 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 245 typename TypeTraits<T>::ParameterType b) const 246 { 247 return a | b; 248 } bit_orcv::cuda::device::bit_or249 __host__ __device__ __forceinline__ bit_or() {} bit_orcv::cuda::device::bit_or250 __host__ __device__ __forceinline__ bit_or(const bit_or&) {} 251 }; 252 253 template <typename T> struct bit_xor : binary_function<T, T, T> 254 { operator ()cv::cuda::device::bit_xor255 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a, 256 typename TypeTraits<T>::ParameterType b) const 257 { 258 return a ^ b; 259 } bit_xorcv::cuda::device::bit_xor260 __host__ __device__ __forceinline__ bit_xor() {} bit_xorcv::cuda::device::bit_xor261 __host__ __device__ __forceinline__ bit_xor(const bit_xor&) {} 262 }; 263 264 template <typename T> struct bit_not : unary_function<T, T> 265 { operator ()cv::cuda::device::bit_not266 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType v) const 267 { 268 return ~v; 269 } bit_notcv::cuda::device::bit_not270 __host__ __device__ __forceinline__ bit_not() {} bit_notcv::cuda::device::bit_not271 __host__ __device__ __forceinline__ bit_not(const bit_not&) {} 272 }; 273 274 // Generalized Identity Operations 275 template <typename T> struct identity : unary_function<T, T> 276 { operator ()cv::cuda::device::identity277 __device__ __forceinline__ typename TypeTraits<T>::ParameterType operator()(typename TypeTraits<T>::ParameterType x) const 278 { 279 return x; 280 } identitycv::cuda::device::identity281 __host__ __device__ __forceinline__ identity() {} identitycv::cuda::device::identity282 __host__ __device__ __forceinline__ identity(const identity&) {} 283 }; 284 285 template <typename T1, typename T2> struct project1st : binary_function<T1, T2, T1> 286 { operator ()cv::cuda::device::project1st287 __device__ __forceinline__ typename TypeTraits<T1>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const 288 { 289 return lhs; 290 } project1stcv::cuda::device::project1st291 __host__ __device__ __forceinline__ project1st() {} project1stcv::cuda::device::project1st292 __host__ __device__ __forceinline__ project1st(const project1st&) {} 293 }; 294 295 template <typename T1, typename T2> struct project2nd : binary_function<T1, T2, T2> 296 { operator ()cv::cuda::device::project2nd297 __device__ __forceinline__ typename TypeTraits<T2>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const 298 { 299 return rhs; 300 } project2ndcv::cuda::device::project2nd301 __host__ __device__ __forceinline__ project2nd() {} project2ndcv::cuda::device::project2nd302 __host__ __device__ __forceinline__ project2nd(const project2nd&) {} 303 }; 304 305 // Min/Max Operations 306 307 #define OPENCV_CUDA_IMPLEMENT_MINMAX(name, type, op) \ 308 template <> struct name<type> : binary_function<type, type, type> \ 309 { \ 310 __device__ __forceinline__ type operator()(type lhs, type rhs) const {return op(lhs, rhs);} \ 311 __host__ __device__ __forceinline__ name() {}\ 312 __host__ __device__ __forceinline__ name(const name&) {}\ 313 }; 314 315 template <typename T> struct maximum : binary_function<T, T, T> 316 { operator ()cv::cuda::device::maximum317 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType lhs, typename TypeTraits<T>::ParameterType rhs) const 318 { 319 return max(lhs, rhs); 320 } maximumcv::cuda::device::maximum321 __host__ __device__ __forceinline__ maximum() {} maximumcv::cuda::device::maximum322 __host__ __device__ __forceinline__ maximum(const maximum&) {} 323 }; 324 325 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, uchar, ::max) 326 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, schar, ::max) 327 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, char, ::max) 328 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, ushort, ::max) 329 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, short, ::max) 330 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, int, ::max) 331 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, uint, ::max) 332 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, float, ::fmax) 333 OPENCV_CUDA_IMPLEMENT_MINMAX(maximum, double, ::fmax) 334 335 template <typename T> struct minimum : binary_function<T, T, T> 336 { operator ()cv::cuda::device::minimum337 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType lhs, typename TypeTraits<T>::ParameterType rhs) const 338 { 339 return min(lhs, rhs); 340 } minimumcv::cuda::device::minimum341 __host__ __device__ __forceinline__ minimum() {} minimumcv::cuda::device::minimum342 __host__ __device__ __forceinline__ minimum(const minimum&) {} 343 }; 344 345 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, uchar, ::min) 346 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, schar, ::min) 347 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, char, ::min) 348 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, ushort, ::min) 349 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, short, ::min) 350 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, int, ::min) 351 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, uint, ::min) 352 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, float, ::fmin) 353 OPENCV_CUDA_IMPLEMENT_MINMAX(minimum, double, ::fmin) 354 355 #undef OPENCV_CUDA_IMPLEMENT_MINMAX 356 357 // Math functions 358 359 template <typename T> struct abs_func : unary_function<T, T> 360 { operator ()cv::cuda::device::abs_func361 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType x) const 362 { 363 return abs(x); 364 } 365 abs_funccv::cuda::device::abs_func366 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func367 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 368 }; 369 template <> struct abs_func<unsigned char> : unary_function<unsigned char, unsigned char> 370 { operator ()cv::cuda::device::abs_func371 __device__ __forceinline__ unsigned char operator ()(unsigned char x) const 372 { 373 return x; 374 } 375 abs_funccv::cuda::device::abs_func376 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func377 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 378 }; 379 template <> struct abs_func<signed char> : unary_function<signed char, signed char> 380 { operator ()cv::cuda::device::abs_func381 __device__ __forceinline__ signed char operator ()(signed char x) const 382 { 383 return ::abs((int)x); 384 } 385 abs_funccv::cuda::device::abs_func386 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func387 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 388 }; 389 template <> struct abs_func<char> : unary_function<char, char> 390 { operator ()cv::cuda::device::abs_func391 __device__ __forceinline__ char operator ()(char x) const 392 { 393 return ::abs((int)x); 394 } 395 abs_funccv::cuda::device::abs_func396 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func397 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 398 }; 399 template <> struct abs_func<unsigned short> : unary_function<unsigned short, unsigned short> 400 { operator ()cv::cuda::device::abs_func401 __device__ __forceinline__ unsigned short operator ()(unsigned short x) const 402 { 403 return x; 404 } 405 abs_funccv::cuda::device::abs_func406 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func407 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 408 }; 409 template <> struct abs_func<short> : unary_function<short, short> 410 { operator ()cv::cuda::device::abs_func411 __device__ __forceinline__ short operator ()(short x) const 412 { 413 return ::abs((int)x); 414 } 415 abs_funccv::cuda::device::abs_func416 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func417 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 418 }; 419 template <> struct abs_func<unsigned int> : unary_function<unsigned int, unsigned int> 420 { operator ()cv::cuda::device::abs_func421 __device__ __forceinline__ unsigned int operator ()(unsigned int x) const 422 { 423 return x; 424 } 425 abs_funccv::cuda::device::abs_func426 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func427 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 428 }; 429 template <> struct abs_func<int> : unary_function<int, int> 430 { operator ()cv::cuda::device::abs_func431 __device__ __forceinline__ int operator ()(int x) const 432 { 433 return ::abs(x); 434 } 435 abs_funccv::cuda::device::abs_func436 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func437 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 438 }; 439 template <> struct abs_func<float> : unary_function<float, float> 440 { operator ()cv::cuda::device::abs_func441 __device__ __forceinline__ float operator ()(float x) const 442 { 443 return ::fabsf(x); 444 } 445 abs_funccv::cuda::device::abs_func446 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func447 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 448 }; 449 template <> struct abs_func<double> : unary_function<double, double> 450 { operator ()cv::cuda::device::abs_func451 __device__ __forceinline__ double operator ()(double x) const 452 { 453 return ::fabs(x); 454 } 455 abs_funccv::cuda::device::abs_func456 __host__ __device__ __forceinline__ abs_func() {} abs_funccv::cuda::device::abs_func457 __host__ __device__ __forceinline__ abs_func(const abs_func&) {} 458 }; 459 460 #define OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(name, func) \ 461 template <typename T> struct name ## _func : unary_function<T, float> \ 462 { \ 463 __device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v) const \ 464 { \ 465 return func ## f(v); \ 466 } \ 467 __host__ __device__ __forceinline__ name ## _func() {} \ 468 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 469 }; \ 470 template <> struct name ## _func<double> : unary_function<double, double> \ 471 { \ 472 __device__ __forceinline__ double operator ()(double v) const \ 473 { \ 474 return func(v); \ 475 } \ 476 __host__ __device__ __forceinline__ name ## _func() {} \ 477 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 478 }; 479 480 #define OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(name, func) \ 481 template <typename T> struct name ## _func : binary_function<T, T, float> \ 482 { \ 483 __device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v1, typename TypeTraits<T>::ParameterType v2) const \ 484 { \ 485 return func ## f(v1, v2); \ 486 } \ 487 __host__ __device__ __forceinline__ name ## _func() {} \ 488 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 489 }; \ 490 template <> struct name ## _func<double> : binary_function<double, double, double> \ 491 { \ 492 __device__ __forceinline__ double operator ()(double v1, double v2) const \ 493 { \ 494 return func(v1, v2); \ 495 } \ 496 __host__ __device__ __forceinline__ name ## _func() {} \ 497 __host__ __device__ __forceinline__ name ## _func(const name ## _func&) {} \ 498 }; 499 500 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sqrt, ::sqrt) 501 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp, ::exp) 502 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp2, ::exp2) 503 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(exp10, ::exp10) 504 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log, ::log) 505 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log2, ::log2) 506 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(log10, ::log10) 507 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sin, ::sin) 508 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(cos, ::cos) 509 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(tan, ::tan) 510 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(asin, ::asin) 511 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(acos, ::acos) 512 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(atan, ::atan) 513 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(sinh, ::sinh) 514 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(cosh, ::cosh) 515 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(tanh, ::tanh) 516 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(asinh, ::asinh) 517 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(acosh, ::acosh) 518 OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR(atanh, ::atanh) 519 520 OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(hypot, ::hypot) 521 OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(atan2, ::atan2) 522 OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR(pow, ::pow) 523 524 #undef OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR 525 #undef OPENCV_CUDA_IMPLEMENT_UN_FUNCTOR_NO_DOUBLE 526 #undef OPENCV_CUDA_IMPLEMENT_BIN_FUNCTOR 527 528 template<typename T> struct hypot_sqr_func : binary_function<T, T, float> 529 { operator ()cv::cuda::device::hypot_sqr_func530 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType src1, typename TypeTraits<T>::ParameterType src2) const 531 { 532 return src1 * src1 + src2 * src2; 533 } hypot_sqr_funccv::cuda::device::hypot_sqr_func534 __host__ __device__ __forceinline__ hypot_sqr_func() {} hypot_sqr_funccv::cuda::device::hypot_sqr_func535 __host__ __device__ __forceinline__ hypot_sqr_func(const hypot_sqr_func&) {} 536 }; 537 538 // Saturate Cast Functor 539 template <typename T, typename D> struct saturate_cast_func : unary_function<T, D> 540 { operator ()cv::cuda::device::saturate_cast_func541 __device__ __forceinline__ D operator ()(typename TypeTraits<T>::ParameterType v) const 542 { 543 return saturate_cast<D>(v); 544 } saturate_cast_funccv::cuda::device::saturate_cast_func545 __host__ __device__ __forceinline__ saturate_cast_func() {} saturate_cast_funccv::cuda::device::saturate_cast_func546 __host__ __device__ __forceinline__ saturate_cast_func(const saturate_cast_func&) {} 547 }; 548 549 // Threshold Functors 550 template <typename T> struct thresh_binary_func : unary_function<T, T> 551 { thresh_binary_funccv::cuda::device::thresh_binary_func552 __host__ __device__ __forceinline__ thresh_binary_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {} 553 operator ()cv::cuda::device::thresh_binary_func554 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 555 { 556 return (src > thresh) * maxVal; 557 } 558 thresh_binary_funccv::cuda::device::thresh_binary_func559 __host__ __device__ __forceinline__ thresh_binary_func() {} thresh_binary_funccv::cuda::device::thresh_binary_func560 __host__ __device__ __forceinline__ thresh_binary_func(const thresh_binary_func& other) 561 : thresh(other.thresh), maxVal(other.maxVal) {} 562 563 T thresh; 564 T maxVal; 565 }; 566 567 template <typename T> struct thresh_binary_inv_func : unary_function<T, T> 568 { thresh_binary_inv_funccv::cuda::device::thresh_binary_inv_func569 __host__ __device__ __forceinline__ thresh_binary_inv_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {} 570 operator ()cv::cuda::device::thresh_binary_inv_func571 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 572 { 573 return (src <= thresh) * maxVal; 574 } 575 thresh_binary_inv_funccv::cuda::device::thresh_binary_inv_func576 __host__ __device__ __forceinline__ thresh_binary_inv_func() {} thresh_binary_inv_funccv::cuda::device::thresh_binary_inv_func577 __host__ __device__ __forceinline__ thresh_binary_inv_func(const thresh_binary_inv_func& other) 578 : thresh(other.thresh), maxVal(other.maxVal) {} 579 580 T thresh; 581 T maxVal; 582 }; 583 584 template <typename T> struct thresh_trunc_func : unary_function<T, T> 585 { thresh_trunc_funccv::cuda::device::thresh_trunc_func586 explicit __host__ __device__ __forceinline__ thresh_trunc_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;} 587 operator ()cv::cuda::device::thresh_trunc_func588 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 589 { 590 return minimum<T>()(src, thresh); 591 } 592 thresh_trunc_funccv::cuda::device::thresh_trunc_func593 __host__ __device__ __forceinline__ thresh_trunc_func() {} thresh_trunc_funccv::cuda::device::thresh_trunc_func594 __host__ __device__ __forceinline__ thresh_trunc_func(const thresh_trunc_func& other) 595 : thresh(other.thresh) {} 596 597 T thresh; 598 }; 599 600 template <typename T> struct thresh_to_zero_func : unary_function<T, T> 601 { thresh_to_zero_funccv::cuda::device::thresh_to_zero_func602 explicit __host__ __device__ __forceinline__ thresh_to_zero_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;} 603 operator ()cv::cuda::device::thresh_to_zero_func604 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 605 { 606 return (src > thresh) * src; 607 } 608 thresh_to_zero_funccv::cuda::device::thresh_to_zero_func609 __host__ __device__ __forceinline__ thresh_to_zero_func() {} thresh_to_zero_funccv::cuda::device::thresh_to_zero_func610 __host__ __device__ __forceinline__ thresh_to_zero_func(const thresh_to_zero_func& other) 611 : thresh(other.thresh) {} 612 613 T thresh; 614 }; 615 616 template <typename T> struct thresh_to_zero_inv_func : unary_function<T, T> 617 { thresh_to_zero_inv_funccv::cuda::device::thresh_to_zero_inv_func618 explicit __host__ __device__ __forceinline__ thresh_to_zero_inv_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;} 619 operator ()cv::cuda::device::thresh_to_zero_inv_func620 __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const 621 { 622 return (src <= thresh) * src; 623 } 624 thresh_to_zero_inv_funccv::cuda::device::thresh_to_zero_inv_func625 __host__ __device__ __forceinline__ thresh_to_zero_inv_func() {} thresh_to_zero_inv_funccv::cuda::device::thresh_to_zero_inv_func626 __host__ __device__ __forceinline__ thresh_to_zero_inv_func(const thresh_to_zero_inv_func& other) 627 : thresh(other.thresh) {} 628 629 T thresh; 630 }; 631 632 // Function Object Adaptors 633 template <typename Predicate> struct unary_negate : unary_function<typename Predicate::argument_type, bool> 634 { unary_negatecv::cuda::device::unary_negate635 explicit __host__ __device__ __forceinline__ unary_negate(const Predicate& p) : pred(p) {} 636 operator ()cv::cuda::device::unary_negate637 __device__ __forceinline__ bool operator()(typename TypeTraits<typename Predicate::argument_type>::ParameterType x) const 638 { 639 return !pred(x); 640 } 641 unary_negatecv::cuda::device::unary_negate642 __host__ __device__ __forceinline__ unary_negate() {} unary_negatecv::cuda::device::unary_negate643 __host__ __device__ __forceinline__ unary_negate(const unary_negate& other) : pred(other.pred) {} 644 645 Predicate pred; 646 }; 647 not1(const Predicate & pred)648 template <typename Predicate> __host__ __device__ __forceinline__ unary_negate<Predicate> not1(const Predicate& pred) 649 { 650 return unary_negate<Predicate>(pred); 651 } 652 653 template <typename Predicate> struct binary_negate : binary_function<typename Predicate::first_argument_type, typename Predicate::second_argument_type, bool> 654 { binary_negatecv::cuda::device::binary_negate655 explicit __host__ __device__ __forceinline__ binary_negate(const Predicate& p) : pred(p) {} 656 operator ()cv::cuda::device::binary_negate657 __device__ __forceinline__ bool operator()(typename TypeTraits<typename Predicate::first_argument_type>::ParameterType x, 658 typename TypeTraits<typename Predicate::second_argument_type>::ParameterType y) const 659 { 660 return !pred(x,y); 661 } 662 binary_negatecv::cuda::device::binary_negate663 __host__ __device__ __forceinline__ binary_negate() {} binary_negatecv::cuda::device::binary_negate664 __host__ __device__ __forceinline__ binary_negate(const binary_negate& other) : pred(other.pred) {} 665 666 Predicate pred; 667 }; 668 not2(const BinaryPredicate & pred)669 template <typename BinaryPredicate> __host__ __device__ __forceinline__ binary_negate<BinaryPredicate> not2(const BinaryPredicate& pred) 670 { 671 return binary_negate<BinaryPredicate>(pred); 672 } 673 674 template <typename Op> struct binder1st : unary_function<typename Op::second_argument_type, typename Op::result_type> 675 { binder1stcv::cuda::device::binder1st676 __host__ __device__ __forceinline__ binder1st(const Op& op_, const typename Op::first_argument_type& arg1_) : op(op_), arg1(arg1_) {} 677 operator ()cv::cuda::device::binder1st678 __device__ __forceinline__ typename Op::result_type operator ()(typename TypeTraits<typename Op::second_argument_type>::ParameterType a) const 679 { 680 return op(arg1, a); 681 } 682 binder1stcv::cuda::device::binder1st683 __host__ __device__ __forceinline__ binder1st() {} binder1stcv::cuda::device::binder1st684 __host__ __device__ __forceinline__ binder1st(const binder1st& other) : op(other.op), arg1(other.arg1) {} 685 686 Op op; 687 typename Op::first_argument_type arg1; 688 }; 689 bind1st(const Op & op,const T & x)690 template <typename Op, typename T> __host__ __device__ __forceinline__ binder1st<Op> bind1st(const Op& op, const T& x) 691 { 692 return binder1st<Op>(op, typename Op::first_argument_type(x)); 693 } 694 695 template <typename Op> struct binder2nd : unary_function<typename Op::first_argument_type, typename Op::result_type> 696 { binder2ndcv::cuda::device::binder2nd697 __host__ __device__ __forceinline__ binder2nd(const Op& op_, const typename Op::second_argument_type& arg2_) : op(op_), arg2(arg2_) {} 698 operator ()cv::cuda::device::binder2nd699 __forceinline__ __device__ typename Op::result_type operator ()(typename TypeTraits<typename Op::first_argument_type>::ParameterType a) const 700 { 701 return op(a, arg2); 702 } 703 binder2ndcv::cuda::device::binder2nd704 __host__ __device__ __forceinline__ binder2nd() {} binder2ndcv::cuda::device::binder2nd705 __host__ __device__ __forceinline__ binder2nd(const binder2nd& other) : op(other.op), arg2(other.arg2) {} 706 707 Op op; 708 typename Op::second_argument_type arg2; 709 }; 710 bind2nd(const Op & op,const T & x)711 template <typename Op, typename T> __host__ __device__ __forceinline__ binder2nd<Op> bind2nd(const Op& op, const T& x) 712 { 713 return binder2nd<Op>(op, typename Op::second_argument_type(x)); 714 } 715 716 // Functor Traits 717 template <typename F> struct IsUnaryFunction 718 { 719 typedef char Yes; 720 struct No {Yes a[2];}; 721 722 template <typename T, typename D> static Yes check(unary_function<T, D>); 723 static No check(...); 724 725 static F makeF(); 726 727 enum { value = (sizeof(check(makeF())) == sizeof(Yes)) }; 728 }; 729 730 template <typename F> struct IsBinaryFunction 731 { 732 typedef char Yes; 733 struct No {Yes a[2];}; 734 735 template <typename T1, typename T2, typename D> static Yes check(binary_function<T1, T2, D>); 736 static No check(...); 737 738 static F makeF(); 739 740 enum { value = (sizeof(check(makeF())) == sizeof(Yes)) }; 741 }; 742 743 namespace functional_detail 744 { 745 template <size_t src_elem_size, size_t dst_elem_size> struct UnOpShift { enum { shift = 1 }; }; 746 template <size_t src_elem_size> struct UnOpShift<src_elem_size, 1> { enum { shift = 4 }; }; 747 template <size_t src_elem_size> struct UnOpShift<src_elem_size, 2> { enum { shift = 2 }; }; 748 749 template <typename T, typename D> struct DefaultUnaryShift 750 { 751 enum { shift = UnOpShift<sizeof(T), sizeof(D)>::shift }; 752 }; 753 754 template <size_t src_elem_size1, size_t src_elem_size2, size_t dst_elem_size> struct BinOpShift { enum { shift = 1 }; }; 755 template <size_t src_elem_size1, size_t src_elem_size2> struct BinOpShift<src_elem_size1, src_elem_size2, 1> { enum { shift = 4 }; }; 756 template <size_t src_elem_size1, size_t src_elem_size2> struct BinOpShift<src_elem_size1, src_elem_size2, 2> { enum { shift = 2 }; }; 757 758 template <typename T1, typename T2, typename D> struct DefaultBinaryShift 759 { 760 enum { shift = BinOpShift<sizeof(T1), sizeof(T2), sizeof(D)>::shift }; 761 }; 762 763 template <typename Func, bool unary = IsUnaryFunction<Func>::value> struct ShiftDispatcher; 764 template <typename Func> struct ShiftDispatcher<Func, true> 765 { 766 enum { shift = DefaultUnaryShift<typename Func::argument_type, typename Func::result_type>::shift }; 767 }; 768 template <typename Func> struct ShiftDispatcher<Func, false> 769 { 770 enum { shift = DefaultBinaryShift<typename Func::first_argument_type, typename Func::second_argument_type, typename Func::result_type>::shift }; 771 }; 772 } 773 774 template <typename Func> struct DefaultTransformShift 775 { 776 enum { shift = functional_detail::ShiftDispatcher<Func>::shift }; 777 }; 778 779 template <typename Func> struct DefaultTransformFunctorTraits 780 { 781 enum { simple_block_dim_x = 16 }; 782 enum { simple_block_dim_y = 16 }; 783 784 enum { smart_block_dim_x = 16 }; 785 enum { smart_block_dim_y = 16 }; 786 enum { smart_shift = DefaultTransformShift<Func>::shift }; 787 }; 788 789 template <typename Func> struct TransformFunctorTraits : DefaultTransformFunctorTraits<Func> {}; 790 791 #define OPENCV_CUDA_TRANSFORM_FUNCTOR_TRAITS(type) \ 792 template <> struct TransformFunctorTraits< type > : DefaultTransformFunctorTraits< type > 793 }}} // namespace cv { namespace cuda { namespace cudev 794 795 //! @endcond 796 797 #endif // __OPENCV_CUDA_FUNCTIONAL_HPP__ 798