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 // Copyright (C) 2013, OpenCV Foundation, all rights reserved. 16 // Third party copyrights are property of their respective owners. 17 // 18 // Redistribution and use in source and binary forms, with or without modification, 19 // are permitted provided that the following conditions are met: 20 // 21 // * Redistribution's of source code must retain the above copyright notice, 22 // this list of conditions and the following disclaimer. 23 // 24 // * Redistribution's in binary form must reproduce the above copyright notice, 25 // this list of conditions and the following disclaimer in the documentation 26 // and/or other materials provided with the distribution. 27 // 28 // * The name of the copyright holders may not be used to endorse or promote products 29 // derived from this software without specific prior written permission. 30 // 31 // This software is provided by the copyright holders and contributors "as is" and 32 // any express or implied warranties, including, but not limited to, the implied 33 // warranties of merchantability and fitness for a particular purpose are disclaimed. 34 // In no event shall the Intel Corporation or contributors be liable for any direct, 35 // indirect, incidental, special, exemplary, or consequential damages 36 // (including, but not limited to, procurement of substitute goods or services; 37 // loss of use, data, or profits; or business interruption) however caused 38 // and on any theory of liability, whether in contract, strict liability, 39 // or tort (including negligence or otherwise) arising in any way out of 40 // the use of this software, even if advised of the possibility of such damage. 41 // 42 //M*/ 43 44 #pragma once 45 46 #ifndef __OPENCV_CUDEV_FUNCTIONAL_FUNCTIONAL_HPP__ 47 #define __OPENCV_CUDEV_FUNCTIONAL_FUNCTIONAL_HPP__ 48 49 #include "../common.hpp" 50 #include "../util/saturate_cast.hpp" 51 #include "../util/vec_traits.hpp" 52 #include "../util/vec_math.hpp" 53 #include "../util/type_traits.hpp" 54 55 namespace cv { namespace cudev { 56 57 //! @addtogroup cudev 58 //! @{ 59 60 // Function Objects 61 62 template <typename _Arg, typename _Result> struct unary_function 63 { 64 typedef _Arg argument_type; 65 typedef _Result result_type; 66 }; 67 68 template <typename _Arg1, typename _Arg2, typename _Result> struct binary_function 69 { 70 typedef _Arg1 first_argument_type; 71 typedef _Arg2 second_argument_type; 72 typedef _Result result_type; 73 }; 74 75 // Arithmetic Operations 76 77 template <typename T> struct plus : binary_function<T, T, T> 78 { 79 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 80 typename TypeTraits<T>::parameter_type b) const 81 { 82 return saturate_cast<T>(a + b); 83 } 84 }; 85 86 template <typename T> struct minus : binary_function<T, T, T> 87 { 88 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 89 typename TypeTraits<T>::parameter_type b) const 90 { 91 return saturate_cast<T>(a - b); 92 } 93 }; 94 95 template <typename T> struct multiplies : binary_function<T, T, T> 96 { 97 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 98 typename TypeTraits<T>::parameter_type b) const 99 { 100 return saturate_cast<T>(a * b); 101 } 102 }; 103 104 template <typename T> struct divides : binary_function<T, T, T> 105 { 106 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 107 typename TypeTraits<T>::parameter_type b) const 108 { 109 return saturate_cast<T>(a / b); 110 } 111 }; 112 113 template <typename T> struct modulus : binary_function<T, T, T> 114 { 115 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 116 typename TypeTraits<T>::parameter_type b) const 117 { 118 return saturate_cast<T>(a % b); 119 } 120 }; 121 122 template <typename T> struct negate : unary_function<T, T> 123 { 124 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a) const 125 { 126 return saturate_cast<T>(-a); 127 } 128 }; 129 130 // Comparison Operations 131 132 template <typename T> struct equal_to : binary_function<T, T, typename MakeVec<uchar, VecTraits<T>::cn>::type> 133 { 134 __device__ __forceinline__ typename MakeVec<uchar, VecTraits<T>::cn>::type 135 operator ()(typename TypeTraits<T>::parameter_type a, 136 typename TypeTraits<T>::parameter_type b) const 137 { 138 return a == b; 139 } 140 }; 141 142 template <typename T> struct not_equal_to : binary_function<T, T, typename MakeVec<uchar, VecTraits<T>::cn>::type> 143 { 144 __device__ __forceinline__ typename MakeVec<uchar, VecTraits<T>::cn>::type 145 operator ()(typename TypeTraits<T>::parameter_type a, 146 typename TypeTraits<T>::parameter_type b) const 147 { 148 return a != b; 149 } 150 }; 151 152 template <typename T> struct greater : binary_function<T, T, typename MakeVec<uchar, VecTraits<T>::cn>::type> 153 { 154 __device__ __forceinline__ typename MakeVec<uchar, VecTraits<T>::cn>::type 155 operator ()(typename TypeTraits<T>::parameter_type a, 156 typename TypeTraits<T>::parameter_type b) const 157 { 158 return a > b; 159 } 160 }; 161 162 template <typename T> struct less : binary_function<T, T, typename MakeVec<uchar, VecTraits<T>::cn>::type> 163 { 164 __device__ __forceinline__ typename MakeVec<uchar, VecTraits<T>::cn>::type 165 operator ()(typename TypeTraits<T>::parameter_type a, 166 typename TypeTraits<T>::parameter_type b) const 167 { 168 return a < b; 169 } 170 }; 171 172 template <typename T> struct greater_equal : binary_function<T, T, typename MakeVec<uchar, VecTraits<T>::cn>::type> 173 { 174 __device__ __forceinline__ typename MakeVec<uchar, VecTraits<T>::cn>::type 175 operator ()(typename TypeTraits<T>::parameter_type a, 176 typename TypeTraits<T>::parameter_type b) const 177 { 178 return a >= b; 179 } 180 }; 181 182 template <typename T> struct less_equal : binary_function<T, T, typename MakeVec<uchar, VecTraits<T>::cn>::type> 183 { 184 __device__ __forceinline__ typename MakeVec<uchar, VecTraits<T>::cn>::type 185 operator ()(typename TypeTraits<T>::parameter_type a, 186 typename TypeTraits<T>::parameter_type b) const 187 { 188 return a <= b; 189 } 190 }; 191 192 // Logical Operations 193 194 template <typename T> struct logical_and : binary_function<T, T, typename MakeVec<uchar, VecTraits<T>::cn>::type> 195 { 196 __device__ __forceinline__ typename MakeVec<uchar, VecTraits<T>::cn>::type 197 operator ()(typename TypeTraits<T>::parameter_type a, 198 typename TypeTraits<T>::parameter_type b) const 199 { 200 return a && b; 201 } 202 }; 203 204 template <typename T> struct logical_or : binary_function<T, T, typename MakeVec<uchar, VecTraits<T>::cn>::type> 205 { 206 __device__ __forceinline__ typename MakeVec<uchar, VecTraits<T>::cn>::type 207 operator ()(typename TypeTraits<T>::parameter_type a, 208 typename TypeTraits<T>::parameter_type b) const 209 { 210 return a || b; 211 } 212 }; 213 214 template <typename T> struct logical_not : unary_function<T, typename MakeVec<uchar, VecTraits<T>::cn>::type> 215 { 216 __device__ __forceinline__ typename MakeVec<uchar, VecTraits<T>::cn>::type 217 operator ()(typename TypeTraits<T>::parameter_type a) const 218 { 219 return !a; 220 } 221 }; 222 223 // Bitwise Operations 224 225 template <typename T> struct bit_and : binary_function<T, T, T> 226 { 227 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 228 typename TypeTraits<T>::parameter_type b) const 229 { 230 return a & b; 231 } 232 }; 233 234 template <typename T> struct bit_or : binary_function<T, T, T> 235 { 236 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 237 typename TypeTraits<T>::parameter_type b) const 238 { 239 return a | b; 240 } 241 }; 242 243 template <typename T> struct bit_xor : binary_function<T, T, T> 244 { 245 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 246 typename TypeTraits<T>::parameter_type b) const 247 { 248 return a ^ b; 249 } 250 }; 251 252 template <typename T> struct bit_not : unary_function<T, T> 253 { 254 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type v) const 255 { 256 return ~v; 257 } 258 }; 259 260 template <typename T> struct bit_lshift : binary_function<T, T, T> 261 { 262 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 263 typename TypeTraits<T>::parameter_type b) const 264 { 265 return a << b; 266 } 267 }; 268 269 template <typename T> struct bit_rshift : binary_function<T, T, T> 270 { 271 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 272 typename TypeTraits<T>::parameter_type b) const 273 { 274 return a >> b; 275 } 276 }; 277 278 // Generalized Identity Operations 279 280 template <typename T> struct identity : unary_function<T, T> 281 { 282 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type x) const 283 { 284 return x; 285 } 286 }; 287 288 template <typename T1, typename T2> struct project1st : binary_function<T1, T2, T1> 289 { 290 __device__ __forceinline__ T1 291 operator ()(typename TypeTraits<T1>::parameter_type lhs, 292 typename TypeTraits<T2>::parameter_type) const 293 { 294 return lhs; 295 } 296 }; 297 298 template <typename T1, typename T2> struct project2nd : binary_function<T1, T2, T2> 299 { 300 __device__ __forceinline__ T2 301 operator ()(typename TypeTraits<T1>::parameter_type, 302 typename TypeTraits<T2>::parameter_type rhs) const 303 { 304 return rhs; 305 } 306 }; 307 308 // Min/Max Operations 309 310 template <typename T> struct maximum : binary_function<T, T, T> 311 { 312 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 313 typename TypeTraits<T>::parameter_type b) const 314 { 315 return max(a, b); 316 } 317 }; 318 319 template <typename T> struct minimum : binary_function<T, T, T> 320 { 321 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, 322 typename TypeTraits<T>::parameter_type b) const 323 { 324 return min(a, b); 325 } 326 }; 327 328 #define CV_CUDEV_MINMAX_INST(type, maxop, minop) \ 329 template <> struct maximum<type> : binary_function<type, type, type> \ 330 { \ 331 __device__ __forceinline__ type operator ()(type a, type b) const {return maxop(a, b);} \ 332 }; \ 333 template <> struct minimum<type> : binary_function<type, type, type> \ 334 { \ 335 __device__ __forceinline__ type operator ()(type a, type b) const {return minop(a, b);} \ 336 }; 337 338 339 CV_CUDEV_MINMAX_INST(uchar, ::max, ::min) 340 CV_CUDEV_MINMAX_INST(schar, ::max, ::min) 341 CV_CUDEV_MINMAX_INST(ushort, ::max, ::min) 342 CV_CUDEV_MINMAX_INST(short, ::max, ::min) 343 CV_CUDEV_MINMAX_INST(int, ::max, ::min) 344 CV_CUDEV_MINMAX_INST(uint, ::max, ::min) 345 CV_CUDEV_MINMAX_INST(float, ::fmaxf, ::fminf) 346 CV_CUDEV_MINMAX_INST(double, ::fmax, ::fmin) 347 348 #undef CV_CUDEV_MINMAX_INST 349 350 // abs_func 351 352 template <typename T> struct abs_func : unary_function<T, T> 353 { 354 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type x) const 355 { 356 return abs(x); 357 } 358 }; 359 360 template <> struct abs_func<uchar> : unary_function<uchar, uchar> 361 { 362 __device__ __forceinline__ uchar operator ()(uchar x) const 363 { 364 return x; 365 } 366 }; 367 368 template <> struct abs_func<schar> : unary_function<schar, schar> 369 { 370 __device__ __forceinline__ schar operator ()(schar x) const 371 { 372 return ::abs((int) x); 373 } 374 }; 375 376 template <> struct abs_func<ushort> : unary_function<ushort, ushort> 377 { 378 __device__ __forceinline__ ushort operator ()(ushort x) const 379 { 380 return x; 381 } 382 }; 383 384 template <> struct abs_func<short> : unary_function<short, short> 385 { 386 __device__ __forceinline__ short operator ()(short x) const 387 { 388 return ::abs((int) x); 389 } 390 }; 391 392 template <> struct abs_func<uint> : unary_function<uint, uint> 393 { 394 __device__ __forceinline__ uint operator ()(uint x) const 395 { 396 return x; 397 } 398 }; 399 400 template <> struct abs_func<int> : unary_function<int, int> 401 { 402 __device__ __forceinline__ int operator ()(int x) const 403 { 404 return ::abs(x); 405 } 406 }; 407 408 template <> struct abs_func<float> : unary_function<float, float> 409 { 410 __device__ __forceinline__ float operator ()(float x) const 411 { 412 return ::fabsf(x); 413 } 414 }; 415 416 template <> struct abs_func<double> : unary_function<double, double> 417 { 418 __device__ __forceinline__ double operator ()(double x) const 419 { 420 return ::fabs(x); 421 } 422 }; 423 424 // absdiff_func 425 426 template <typename T> struct absdiff_func : binary_function<T, T, T> 427 { 428 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type a, typename TypeTraits<T>::parameter_type b) const 429 { 430 abs_func<T> f; 431 return f(a - b); 432 } 433 }; 434 435 // Math functions 436 437 template <typename T> struct sqr_func : unary_function<T, T> 438 { 439 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type x) const 440 { 441 return x * x; 442 } 443 }; 444 445 namespace functional_detail 446 { 447 template <typename T> struct FloatType 448 { 449 typedef typename MakeVec< 450 typename LargerType<float, typename VecTraits<T>::elem_type>::type, 451 VecTraits<T>::cn 452 >::type type; 453 }; 454 } 455 456 #define CV_CUDEV_UNARY_FUNCTION_INST(name, func) \ 457 template <typename T> struct name ## _func : unary_function<T, typename functional_detail::FloatType<T>::type> \ 458 { \ 459 __device__ __forceinline__ typename functional_detail::FloatType<T>::type operator ()(typename TypeTraits<T>::parameter_type a) const \ 460 { \ 461 return name(a); \ 462 } \ 463 }; \ 464 template <> struct name ## _func<uchar> : unary_function<uchar, float> \ 465 { \ 466 __device__ __forceinline__ float operator ()(uchar a) const \ 467 { \ 468 return func ## f(a); \ 469 } \ 470 }; \ 471 template <> struct name ## _func<schar> : unary_function<schar, float> \ 472 { \ 473 __device__ __forceinline__ float operator ()(schar a) const \ 474 { \ 475 return func ## f(a); \ 476 } \ 477 }; \ 478 template <> struct name ## _func<ushort> : unary_function<ushort, float> \ 479 { \ 480 __device__ __forceinline__ float operator ()(ushort a) const \ 481 { \ 482 return func ## f(a); \ 483 } \ 484 }; \ 485 template <> struct name ## _func<short> : unary_function<short, float> \ 486 { \ 487 __device__ __forceinline__ float operator ()(short a) const \ 488 { \ 489 return func ## f(a); \ 490 } \ 491 }; \ 492 template <> struct name ## _func<uint> : unary_function<uint, float> \ 493 { \ 494 __device__ __forceinline__ float operator ()(uint a) const \ 495 { \ 496 return func ## f(a); \ 497 } \ 498 }; \ 499 template <> struct name ## _func<int> : unary_function<int, float> \ 500 { \ 501 __device__ __forceinline__ float operator ()(int a) const \ 502 { \ 503 return func ## f(a); \ 504 } \ 505 }; \ 506 template <> struct name ## _func<float> : unary_function<float, float> \ 507 { \ 508 __device__ __forceinline__ float operator ()(float a) const \ 509 { \ 510 return func ## f(a); \ 511 } \ 512 }; \ 513 template <> struct name ## _func<double> : unary_function<double, double> \ 514 { \ 515 __device__ __forceinline__ double operator ()(double a) const \ 516 { \ 517 return func(a); \ 518 } \ 519 }; 520 521 CV_CUDEV_UNARY_FUNCTION_INST(sqrt, ::sqrt) 522 CV_CUDEV_UNARY_FUNCTION_INST(exp, ::exp) 523 CV_CUDEV_UNARY_FUNCTION_INST(exp2, ::exp2) 524 CV_CUDEV_UNARY_FUNCTION_INST(exp10, ::exp10) 525 CV_CUDEV_UNARY_FUNCTION_INST(log, ::log) 526 CV_CUDEV_UNARY_FUNCTION_INST(log2, ::log2) 527 CV_CUDEV_UNARY_FUNCTION_INST(log10, ::log10) 528 CV_CUDEV_UNARY_FUNCTION_INST(sin, ::sin) 529 CV_CUDEV_UNARY_FUNCTION_INST(cos, ::cos) 530 CV_CUDEV_UNARY_FUNCTION_INST(tan, ::tan) 531 CV_CUDEV_UNARY_FUNCTION_INST(asin, ::asin) 532 CV_CUDEV_UNARY_FUNCTION_INST(acos, ::acos) 533 CV_CUDEV_UNARY_FUNCTION_INST(atan, ::atan) 534 CV_CUDEV_UNARY_FUNCTION_INST(sinh, ::sinh) 535 CV_CUDEV_UNARY_FUNCTION_INST(cosh, ::cosh) 536 CV_CUDEV_UNARY_FUNCTION_INST(tanh, ::tanh) 537 CV_CUDEV_UNARY_FUNCTION_INST(asinh, ::asinh) 538 CV_CUDEV_UNARY_FUNCTION_INST(acosh, ::acosh) 539 CV_CUDEV_UNARY_FUNCTION_INST(atanh, ::atanh) 540 541 #undef CV_CUDEV_UNARY_FUNCTION_INST 542 543 #define CV_CUDEV_BINARY_FUNCTION_INST(name, func) \ 544 template <typename T> struct name ## _func : binary_function<T, T, typename functional_detail::FloatType<T>::type> \ 545 { \ 546 __device__ __forceinline__ typename functional_detail::FloatType<T>::type operator ()(typename TypeTraits<T>::parameter_type a, typename TypeTraits<T>::parameter_type b) const \ 547 { \ 548 return name(a, b); \ 549 } \ 550 }; \ 551 template <> struct name ## _func<uchar> : binary_function<uchar, uchar, float> \ 552 { \ 553 __device__ __forceinline__ float operator ()(uchar a, uchar b) const \ 554 { \ 555 return func ## f(a, b); \ 556 } \ 557 }; \ 558 template <> struct name ## _func<schar> : binary_function<schar, schar, float> \ 559 { \ 560 __device__ __forceinline__ float operator ()(schar a, schar b) const \ 561 { \ 562 return func ## f(a, b); \ 563 } \ 564 }; \ 565 template <> struct name ## _func<ushort> : binary_function<ushort, ushort, float> \ 566 { \ 567 __device__ __forceinline__ float operator ()(ushort a, ushort b) const \ 568 { \ 569 return func ## f(a, b); \ 570 } \ 571 }; \ 572 template <> struct name ## _func<short> : binary_function<short, short, float> \ 573 { \ 574 __device__ __forceinline__ float operator ()(short a, short b) const \ 575 { \ 576 return func ## f(a, b); \ 577 } \ 578 }; \ 579 template <> struct name ## _func<uint> : binary_function<uint, uint, float> \ 580 { \ 581 __device__ __forceinline__ float operator ()(uint a, uint b) const \ 582 { \ 583 return func ## f(a, b); \ 584 } \ 585 }; \ 586 template <> struct name ## _func<int> : binary_function<int, int, float> \ 587 { \ 588 __device__ __forceinline__ float operator ()(int a, int b) const \ 589 { \ 590 return func ## f(a, b); \ 591 } \ 592 }; \ 593 template <> struct name ## _func<float> : binary_function<float, float, float> \ 594 { \ 595 __device__ __forceinline__ float operator ()(float a, float b) const \ 596 { \ 597 return func ## f(a, b); \ 598 } \ 599 }; \ 600 template <> struct name ## _func<double> : binary_function<double, double, double> \ 601 { \ 602 __device__ __forceinline__ double operator ()(double a, double b) const \ 603 { \ 604 return func(a, b); \ 605 } \ 606 }; 607 608 CV_CUDEV_BINARY_FUNCTION_INST(hypot, ::hypot) 609 CV_CUDEV_BINARY_FUNCTION_INST(atan2, ::atan2) 610 611 #undef CV_CUDEV_BINARY_FUNCTION_INST 612 613 template <typename T> struct magnitude_func : binary_function<T, T, typename functional_detail::FloatType<T>::type> 614 { 615 __device__ __forceinline__ typename functional_detail::FloatType<T>::type operator ()(typename TypeTraits<T>::parameter_type a, typename TypeTraits<T>::parameter_type b) const 616 { 617 sqrt_func<typename functional_detail::FloatType<T>::type> f; 618 return f(a * a + b * b); 619 } 620 }; 621 622 template <typename T> struct magnitude_sqr_func : binary_function<T, T, typename functional_detail::FloatType<T>::type> 623 { 624 __device__ __forceinline__ typename functional_detail::FloatType<T>::type operator ()(typename TypeTraits<T>::parameter_type a, typename TypeTraits<T>::parameter_type b) const 625 { 626 return a * a + b * b; 627 } 628 }; 629 630 template <typename T, bool angleInDegrees> struct direction_func : binary_function<T, T, T> 631 { 632 __device__ T operator ()(T x, T y) const 633 { 634 atan2_func<T> f; 635 typename atan2_func<T>::result_type angle = f(y, x); 636 637 angle += (angle < 0) * (2.0f * CV_PI_F); 638 639 if (angleInDegrees) 640 angle *= (180.0f / CV_PI_F); 641 642 return saturate_cast<T>(angle); 643 } 644 }; 645 646 template <typename T> struct pow_func : binary_function<T, float, float> 647 { 648 __device__ __forceinline__ float operator ()(T val, float power) const 649 { 650 return ::powf(val, power); 651 } 652 }; 653 template <> struct pow_func<double> : binary_function<double, double, double> 654 { 655 __device__ __forceinline__ double operator ()(double val, double power) const 656 { 657 return ::pow(val, power); 658 } 659 }; 660 661 // Saturate Cast Functor 662 663 template <typename T, typename D> struct saturate_cast_func : unary_function<T, D> 664 { 665 __device__ __forceinline__ D operator ()(typename TypeTraits<T>::parameter_type v) const 666 { 667 return saturate_cast<D>(v); 668 } 669 }; 670 671 // Threshold Functors 672 673 template <typename T> struct ThreshBinaryFunc : unary_function<T, T> 674 { 675 T thresh; 676 T maxVal; 677 678 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type src) const 679 { 680 return saturate_cast<T>(src > thresh) * maxVal; 681 } 682 }; 683 684 template <typename T> 685 __host__ __device__ ThreshBinaryFunc<T> thresh_binary_func(T thresh, T maxVal) 686 { 687 ThreshBinaryFunc<T> f; 688 f.thresh = thresh; 689 f.maxVal = maxVal; 690 return f; 691 } 692 693 template <typename T> struct ThreshBinaryInvFunc : unary_function<T, T> 694 { 695 T thresh; 696 T maxVal; 697 698 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type src) const 699 { 700 return saturate_cast<T>(src <= thresh) * maxVal; 701 } 702 }; 703 704 template <typename T> 705 __host__ __device__ ThreshBinaryInvFunc<T> thresh_binary_inv_func(T thresh, T maxVal) 706 { 707 ThreshBinaryInvFunc<T> f; 708 f.thresh = thresh; 709 f.maxVal = maxVal; 710 return f; 711 } 712 713 template <typename T> struct ThreshTruncFunc : unary_function<T, T> 714 { 715 T thresh; 716 717 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type src) const 718 { 719 minimum<T> minOp; 720 return minOp(src, thresh); 721 } 722 }; 723 724 template <typename T> 725 __host__ __device__ ThreshTruncFunc<T> thresh_trunc_func(T thresh) 726 { 727 ThreshTruncFunc<T> f; 728 f.thresh = thresh; 729 return f; 730 } 731 732 template <typename T> struct ThreshToZeroFunc : unary_function<T, T> 733 { 734 T thresh; 735 736 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type src) const 737 { 738 return saturate_cast<T>(src > thresh) * src; 739 } 740 }; 741 742 template <typename T> 743 __host__ __device__ ThreshToZeroFunc<T> thresh_to_zero_func(T thresh) 744 { 745 ThreshToZeroFunc<T> f; 746 f.thresh = thresh; 747 return f; 748 } 749 750 template <typename T> struct ThreshToZeroInvFunc : unary_function<T, T> 751 { 752 T thresh; 753 754 __device__ __forceinline__ T operator ()(typename TypeTraits<T>::parameter_type src) const 755 { 756 return saturate_cast<T>(src <= thresh) * src; 757 } 758 }; 759 760 template <typename T> 761 __host__ __device__ ThreshToZeroInvFunc<T> thresh_to_zero_inv_func(T thresh) 762 { 763 ThreshToZeroInvFunc<T> f; 764 f.thresh = thresh; 765 return f; 766 } 767 768 // Function Object Adaptors 769 770 template <class Predicate> struct UnaryNegate : unary_function<typename Predicate::argument_type, typename Predicate::result_type> 771 { 772 Predicate pred; 773 774 __device__ __forceinline__ typename Predicate::result_type operator ()( 775 typename TypeTraits<typename Predicate::argument_type>::parameter_type x) const 776 { 777 return !pred(x); 778 } 779 }; 780 781 template <class Predicate> 782 __host__ __device__ UnaryNegate<Predicate> not1(const Predicate& pred) 783 { 784 UnaryNegate<Predicate> n; 785 n.pred = pred; 786 return n; 787 } 788 789 template <class Predicate> struct BinaryNegate : binary_function<typename Predicate::first_argument_type, typename Predicate::second_argument_type, typename Predicate::result_type> 790 { 791 Predicate pred; 792 793 __device__ __forceinline__ typename Predicate::result_type operator ()( 794 typename TypeTraits<typename Predicate::first_argument_type>::parameter_type x, 795 typename TypeTraits<typename Predicate::second_argument_type>::parameter_type y) const 796 { 797 return !pred(x, y); 798 } 799 }; 800 801 template <class Predicate> 802 __host__ __device__ BinaryNegate<Predicate> not2(const Predicate& pred) 803 { 804 BinaryNegate<Predicate> n; 805 n.pred = pred; 806 return n; 807 } 808 809 template <class Op> struct Binder1st : unary_function<typename Op::second_argument_type, typename Op::result_type> 810 { 811 Op op; 812 typename Op::first_argument_type arg1; 813 814 __device__ __forceinline__ typename Op::result_type operator ()( 815 typename TypeTraits<typename Op::second_argument_type>::parameter_type a) const 816 { 817 return op(arg1, a); 818 } 819 }; 820 821 template <class Op> 822 __host__ __device__ Binder1st<Op> bind1st(const Op& op, const typename Op::first_argument_type& arg1) 823 { 824 Binder1st<Op> b; 825 b.op = op; 826 b.arg1 = arg1; 827 return b; 828 } 829 830 template <class Op> struct Binder2nd : unary_function<typename Op::first_argument_type, typename Op::result_type> 831 { 832 Op op; 833 typename Op::second_argument_type arg2; 834 835 __device__ __forceinline__ typename Op::result_type operator ()( 836 typename TypeTraits<typename Op::first_argument_type>::parameter_type a) const 837 { 838 return op(a, arg2); 839 } 840 }; 841 842 template <class Op> 843 __host__ __device__ Binder2nd<Op> bind2nd(const Op& op, const typename Op::second_argument_type& arg2) 844 { 845 Binder2nd<Op> b; 846 b.op = op; 847 b.arg2 = arg2; 848 return b; 849 } 850 851 // Functor Traits 852 853 template <typename F> struct IsUnaryFunction 854 { 855 typedef char Yes; 856 struct No {Yes a[2];}; 857 858 template <typename T, typename D> static Yes check(unary_function<T, D>); 859 static No check(...); 860 861 static F makeF(); 862 863 enum { value = (sizeof(check(makeF())) == sizeof(Yes)) }; 864 }; 865 866 template <typename F> struct IsBinaryFunction 867 { 868 typedef char Yes; 869 struct No {Yes a[2];}; 870 871 template <typename T1, typename T2, typename D> static Yes check(binary_function<T1, T2, D>); 872 static No check(...); 873 874 static F makeF(); 875 876 enum { value = (sizeof(check(makeF())) == sizeof(Yes)) }; 877 }; 878 879 //! @} 880 881 }} 882 883 #endif 884
