modules/gpu/src/opencv2/gpu/device/functional.hpp

   1 /*M///////////////////////////////////////////////////////////////////////////////////////\r
   2 //\r
   3 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
   4 //\r
   5 //  By downloading, copying, installing or using the software you agree to this license.\r
   6 //  If you do not agree to this license, do not download, install,\r
   7 //  copy or use the software.\r
   8 //\r
   9 //\r
  10 //                           License Agreement\r
  11 //                For Open Source Computer Vision Library\r
  12 //\r
  13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.\r
  14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.\r
  15 // Third party copyrights are property of their respective owners.\r
  16 //\r
  17 // Redistribution and use in source and binary forms, with or without modification,\r
  18 // are permitted provided that the following conditions are met:\r
  19 //\r
  20 //   * Redistribution's of source code must retain the above copyright notice,\r
  21 //     this list of conditions and the following disclaimer.\r
  22 //\r
  23 //   * Redistribution's in binary form must reproduce the above copyright notice,\r
  24 //     this list of conditions and the following disclaimer in the documentation\r
  25 //     and/or other materials provided with the distribution.\r
  26 //\r
  27 //   * The name of the copyright holders may not be used to endorse or promote products\r
  28 //     derived from this software without specific prior written permission.\r
  29 //\r
  30 // This software is provided by the copyright holders and contributors "as is" and\r
  31 // any express or implied warranties, including, but not limited to, the implied\r
  32 // warranties of merchantability and fitness for a particular purpose are disclaimed.\r
  33 // In no event shall the Intel Corporation or contributors be liable for any direct,\r
  34 // indirect, incidental, special, exemplary, or consequential damages\r
  35 // (including, but not limited to, procurement of substitute goods or services;\r
  36 // loss of use, data, or profits; or business interruption) however caused\r
  37 // and on any theory of liability, whether in contract, strict liability,\r
  38 // or tort (including negligence or otherwise) arising in any way out of\r
  39 // the use of this software, even if advised of the possibility of such damage.\r
  40 //\r
  41 //M*/\r
  42 \r
  43 #ifndef __OPENCV_GPU_FUNCTIONAL_HPP__\r
  44 #define __OPENCV_GPU_FUNCTIONAL_HPP__\r
  45 \r
  46 #include <thrust/functional.h>\r
  47 #include "saturate_cast.hpp"\r
  48 #include "vec_traits.hpp"\r
  49 #include "type_traits.hpp"\r
  50 \r
  51 namespace cv { namespace gpu { namespace device\r
  52 {\r
  53     // Function Objects\r
  54 \r
  55     using thrust::unary_function;\r
  56     using thrust::binary_function;\r
  57 \r
  58     // Arithmetic Operations\r
  59     template <typename T> struct plus : binary_function<T, T, T>\r
  60     {\r
  61         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
  62                                                  typename TypeTraits<T>::ParameterType b) const\r
  63         {\r
  64             return a + b;\r
  65         }\r
  66         __device__ __forceinline__ plus(const plus& other):binary_function<T,T,T>(){}\r
  67         __device__ __forceinline__ plus():binary_function<T,T,T>(){}\r
  68     };\r
  69 \r
  70     template <typename T> struct minus : binary_function<T, T, T>\r
  71     {\r
  72         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
  73                                                  typename TypeTraits<T>::ParameterType b) const\r
  74         {\r
  75             return a - b;\r
  76         }\r
  77         __device__ __forceinline__ minus(const minus& other):binary_function<T,T,T>(){}\r
  78         __device__ __forceinline__ minus():binary_function<T,T,T>(){}\r
  79     };\r
  80 \r
  81     template <typename T> struct multiplies : binary_function<T, T, T>\r
  82     {\r
  83         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
  84                                                  typename TypeTraits<T>::ParameterType b) const\r
  85         {\r
  86             return a * b;\r
  87         }\r
  88         __device__ __forceinline__ multiplies(const multiplies& other):binary_function<T,T,T>(){}\r
  89         __device__ __forceinline__ multiplies():binary_function<T,T,T>(){}\r
  90     };\r
  91 \r
  92     template <typename T> struct divides : binary_function<T, T, T>\r
  93     {\r
  94         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
  95                                                  typename TypeTraits<T>::ParameterType b) const\r
  96         {\r
  97             return a / b;\r
  98         }\r
  99         __device__ __forceinline__ divides(const divides& other):binary_function<T,T,T>(){}\r
 100         __device__ __forceinline__ divides():binary_function<T,T,T>(){}\r
 101     };\r
 102 \r
 103     template <typename T> struct modulus : binary_function<T, T, T>\r
 104     {\r
 105         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
 106                                                  typename TypeTraits<T>::ParameterType b) const\r
 107         {\r
 108             return a % b;\r
 109         }\r
 110         __device__ __forceinline__ modulus(const modulus& other):binary_function<T,T,T>(){}\r
 111         __device__ __forceinline__ modulus():binary_function<T,T,T>(){}\r
 112     };\r
 113 \r
 114     template <typename T> struct negate : unary_function<T, T>\r
 115     {\r
 116         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a) const\r
 117         {\r
 118             return -a;\r
 119         }\r
 120         __device__ __forceinline__ negate(const negate& other):unary_function<T,T>(){}\r
 121         __device__ __forceinline__ negate():unary_function<T,T>(){}\r
 122     };\r
 123 \r
 124     // Comparison Operations\r
 125     template <typename T> struct equal_to : binary_function<T, T, bool>\r
 126     {\r
 127         __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
 128                                                     typename TypeTraits<T>::ParameterType b) const\r
 129         {\r
 130             return a == b;\r
 131         }\r
 132         __device__ __forceinline__ equal_to(const equal_to& other):binary_function<T,T,bool>(){}\r
 133         __device__ __forceinline__ equal_to():binary_function<T,T,bool>(){}\r
 134     };\r
 135 \r
 136     template <typename T> struct not_equal_to : binary_function<T, T, bool>\r
 137     {\r
 138         __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
 139                                                     typename TypeTraits<T>::ParameterType b) const\r
 140         {\r
 141             return a != b;\r
 142         }\r
 143         __device__ __forceinline__ not_equal_to(const not_equal_to& other):binary_function<T,T,bool>(){}\r
 144         __device__ __forceinline__ not_equal_to():binary_function<T,T,bool>(){}\r
 145     };\r
 146 \r
 147     template <typename T> struct greater : binary_function<T, T, bool>\r
 148     {\r
 149         __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
 150                                                     typename TypeTraits<T>::ParameterType b) const\r
 151         {\r
 152             return a > b;\r
 153         }\r
 154         __device__ __forceinline__ greater(const greater& other):binary_function<T,T,bool>(){}\r
 155         __device__ __forceinline__ greater():binary_function<T,T,bool>(){}\r
 156     };\r
 157 \r
 158     template <typename T> struct less : binary_function<T, T, bool>\r
 159     {\r
 160         __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
 161                                                     typename TypeTraits<T>::ParameterType b) const\r
 162         {\r
 163             return a < b;\r
 164         }\r
 165         __device__ __forceinline__ less(const less& other):binary_function<T,T,bool>(){}\r
 166         __device__ __forceinline__ less():binary_function<T,T,bool>(){}\r
 167     };\r
 168 \r
 169     template <typename T> struct greater_equal : binary_function<T, T, bool>\r
 170     {\r
 171         __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
 172                                                     typename TypeTraits<T>::ParameterType b) const\r
 173         {\r
 174             return a >= b;\r
 175         }\r
 176         __device__ __forceinline__ greater_equal(const greater_equal& other):binary_function<T,T,bool>(){}\r
 177         __device__ __forceinline__ greater_equal():binary_function<T,T,bool>(){}\r
 178     };\r
 179 \r
 180     template <typename T> struct less_equal : binary_function<T, T, bool>\r
 181     {\r
 182         __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
 183                                                     typename TypeTraits<T>::ParameterType b) const\r
 184         {\r
 185             return a <= b;\r
 186         }\r
 187         __device__ __forceinline__ less_equal(const less_equal& other):binary_function<T,T,bool>(){}\r
 188         __device__ __forceinline__ less_equal():binary_function<T,T,bool>(){}\r
 189     };\r
 190 \r
 191     // Logical Operations\r
 192     template <typename T> struct logical_and : binary_function<T, T, bool>\r
 193     {\r
 194         __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
 195                                                     typename TypeTraits<T>::ParameterType b) const\r
 196         {\r
 197             return a && b;\r
 198         }\r
 199         __device__ __forceinline__ logical_and(const logical_and& other):binary_function<T,T,bool>(){}\r
 200         __device__ __forceinline__ logical_and():binary_function<T,T,bool>(){}\r
 201     };\r
 202 \r
 203     template <typename T> struct logical_or : binary_function<T, T, bool>\r
 204     {\r
 205         __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a,\r
 206                                                     typename TypeTraits<T>::ParameterType b) const\r
 207         {\r
 208             return a || b;\r
 209         }\r
 210         __device__ __forceinline__ logical_or(const logical_or& other):binary_function<T,T,bool>(){}\r
 211         __device__ __forceinline__ logical_or():binary_function<T,T,bool>(){}\r
 212     };\r
 213 \r
 214     template <typename T> struct logical_not : unary_function<T, bool>\r
 215     {\r
 216         __device__ __forceinline__ bool operator ()(typename TypeTraits<T>::ParameterType a) const\r
 217         {\r
 218             return !a;\r
 219         }\r
 220         __device__ __forceinline__ logical_not(const logical_not& other):unary_function<T,bool>(){}\r
 221         __device__ __forceinline__ logical_not():unary_function<T,bool>(){}\r
 222     };\r
 223 \r
 224     // Bitwise Operations\r
 225     template <typename T> struct bit_and : binary_function<T, T, T>\r
 226     {\r
 227         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
 228                                                  typename TypeTraits<T>::ParameterType b) const\r
 229         {\r
 230             return a & b;\r
 231         }\r
 232         __device__ __forceinline__ bit_and(const bit_and& other):binary_function<T,T,T>(){}\r
 233         __device__ __forceinline__ bit_and():binary_function<T,T,T>(){}\r
 234     };\r
 235 \r
 236     template <typename T> struct bit_or : binary_function<T, T, T>\r
 237     {\r
 238         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
 239                                                  typename TypeTraits<T>::ParameterType b) const\r
 240         {\r
 241             return a | b;\r
 242         }\r
 243         __device__ __forceinline__ bit_or(const bit_or& other):binary_function<T,T,T>(){}\r
 244         __device__ __forceinline__ bit_or():binary_function<T,T,T>(){}\r
 245     };\r
 246 \r
 247     template <typename T> struct bit_xor : binary_function<T, T, T>\r
 248     {\r
 249         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType a,\r
 250                                                  typename TypeTraits<T>::ParameterType b) const\r
 251         {\r
 252             return a ^ b;\r
 253         }\r
 254         __device__ __forceinline__ bit_xor(const bit_xor& other):binary_function<T,T,T>(){}\r
 255         __device__ __forceinline__ bit_xor():binary_function<T,T,T>(){}\r
 256     };\r
 257 \r
 258     template <typename T> struct bit_not : unary_function<T, T>\r
 259     {\r
 260         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType v) const\r
 261         {\r
 262             return ~v;\r
 263         }\r
 264         __device__ __forceinline__ bit_not(const bit_not& other):unary_function<T,T>(){}\r
 265         __device__ __forceinline__ bit_not():unary_function<T,T>(){}\r
 266     };\r
 267 \r
 268     // Generalized Identity Operations\r
 269     template <typename T> struct identity : unary_function<T, T>\r
 270     {\r
 271         __device__ __forceinline__ typename TypeTraits<T>::ParameterType operator()(typename TypeTraits<T>::ParameterType x) const\r
 272         {\r
 273             return x;\r
 274         }\r
 275         __device__ __forceinline__ identity(const identity& other):unary_function<T,T>(){}\r
 276         __device__ __forceinline__ identity():unary_function<T,T>(){}\r
 277     };\r
 278 \r
 279     template <typename T1, typename T2> struct project1st : binary_function<T1, T2, T1>\r
 280     {\r
 281         __device__ __forceinline__ typename TypeTraits<T1>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const\r
 282         {\r
 283             return lhs;\r
 284         }\r
 285         __device__ __forceinline__ project1st(const project1st& other):binary_function<T1,T2,T1>(){}\r
 286         __device__ __forceinline__ project1st():binary_function<T1,T2,T1>(){}\r
 287     };\r
 288 \r
 289     template <typename T1, typename T2> struct project2nd : binary_function<T1, T2, T2>\r
 290     {\r
 291         __device__ __forceinline__ typename TypeTraits<T2>::ParameterType operator()(typename TypeTraits<T1>::ParameterType lhs, typename TypeTraits<T2>::ParameterType rhs) const\r
 292         {\r
 293             return rhs;\r
 294         }\r
 295         __device__ __forceinline__ project2nd(const project2nd& other):binary_function<T1,T2,T2>(){}\r
 296         __device__ __forceinline__ project2nd():binary_function<T1,T2,T2>(){}\r
 297     };\r
 298 \r
 299     // Min/Max Operations\r
 300 \r
 301 #define OPENCV_GPU_IMPLEMENT_MINMAX(name, type, op) \\r
 302     template <> struct name<type> : binary_function<type, type, type> \\r
 303     { \\r
 304         __device__ __forceinline__ type operator()(type lhs, type rhs) const {return op(lhs, rhs);} \\r
 305         __device__ __forceinline__ name(const name& other):binary_function<type, type, type>(){}\\r
 306         __device__ __forceinline__ name():binary_function<type, type, type>(){}\\r
 307     };\r
 308 \r
 309     template <typename T> struct maximum : binary_function<T, T, T>\r
 310     {\r
 311         __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType lhs, typename TypeTraits<T>::ParameterType rhs) const\r
 312         {\r
 313             return lhs < rhs ? rhs : lhs;\r
 314         }\r
 315         __device__ __forceinline__ maximum(const maximum& other):binary_function<T, T, T>(){}\r
 316         __device__ __forceinline__ maximum():binary_function<T, T, T>(){}\r
 317     };\r
 318 \r
 319     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, uchar, ::max)\r
 320     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, schar, ::max)\r
 321     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, char, ::max)\r
 322     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, ushort, ::max)\r
 323     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, short, ::max)\r
 324     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, int, ::max)\r
 325     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, uint, ::max)\r
 326     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, float, ::fmax)\r
 327     OPENCV_GPU_IMPLEMENT_MINMAX(maximum, double, ::fmax)\r
 328 \r
 329     template <typename T> struct minimum : binary_function<T, T, T>\r
 330     {\r
 331         __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType lhs, typename TypeTraits<T>::ParameterType rhs) const\r
 332         {\r
 333             return lhs < rhs ? lhs : rhs;\r
 334         }\r
 335         __device__ __forceinline__ minimum(const minimum& other):binary_function<T, T, T>(){}\r
 336         __device__ __forceinline__ minimum():binary_function<T, T, T>(){}\r
 337     };\r
 338 \r
 339     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, uchar, ::min)\r
 340     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, schar, ::min)\r
 341     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, char, ::min)\r
 342     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, ushort, ::min)\r
 343     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, short, ::min)\r
 344     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, int, ::min)\r
 345     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, uint, ::min)\r
 346     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, float, ::fmin)\r
 347     OPENCV_GPU_IMPLEMENT_MINMAX(minimum, double, ::fmin)\r
 348 \r
 349 #undef OPENCV_GPU_IMPLEMENT_MINMAX\r
 350 \r
 351     // Math functions\r
 352 ///bound=========================================\r
 353 #define OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(name, func) \\r
 354     template <typename T> struct name ## _func : unary_function<T, float> \\r
 355     { \\r
 356         __device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v) const \\r
 357         { \\r
 358             return func ## f(v); \\r
 359         } \\r
 360     }; \\r
 361     template <> struct name ## _func<double> : unary_function<double, double> \\r
 362     { \\r
 363         __device__ __forceinline__ double operator ()(double v) const \\r
 364         { \\r
 365             return func(v); \\r
 366         } \\r
 367     };\r
 368 \r
 369 #define OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(name, func) \\r
 370     template <typename T> struct name ## _func : binary_function<T, T, float> \\r
 371     { \\r
 372         __device__ __forceinline__ float operator ()(typename TypeTraits<T>::ParameterType v1, typename TypeTraits<T>::ParameterType v2) const \\r
 373         { \\r
 374             return func ## f(v1, v2); \\r
 375         } \\r
 376     }; \\r
 377     template <> struct name ## _func<double> : binary_function<double, double, double> \\r
 378     { \\r
 379         __device__ __forceinline__ double operator ()(double v1, double v2) const \\r
 380         { \\r
 381             return func(v1, v2); \\r
 382         } \\r
 383     };\r
 384 \r
 385     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(fabs, ::fabs)\r
 386     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sqrt, ::sqrt)\r
 387     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp, ::exp)\r
 388     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp2, ::exp2)\r
 389     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(exp10, ::exp10)\r
 390     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log, ::log)\r
 391     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log2, ::log2)\r
 392     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(log10, ::log10)\r
 393     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sin, ::sin)\r
 394     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cos, ::cos)\r
 395     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tan, ::tan)\r
 396     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asin, ::asin)\r
 397     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acos, ::acos)\r
 398     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atan, ::atan)\r
 399     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(sinh, ::sinh)\r
 400     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(cosh, ::cosh)\r
 401     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(tanh, ::tanh)\r
 402     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(asinh, ::asinh)\r
 403     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(acosh, ::acosh)\r
 404     OPENCV_GPU_IMPLEMENT_UN_FUNCTOR(atanh, ::atanh)\r
 405 \r
 406     OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(hypot, ::hypot)\r
 407     OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(atan2, ::atan2)\r
 408     OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR(pow, ::pow)\r
 409 \r
 410     #undef OPENCV_GPU_IMPLEMENT_UN_FUNCTOR\r
 411     #undef OPENCV_GPU_IMPLEMENT_BIN_FUNCTOR\r
 412 \r
 413     template<typename T> struct hypot_sqr_func : binary_function<T, T, float>\r
 414     {\r
 415         __device__ __forceinline__ T operator ()(typename TypeTraits<T>::ParameterType src1, typename TypeTraits<T>::ParameterType src2) const\r
 416         {\r
 417             return src1 * src1 + src2 * src2;\r
 418         }\r
 419         __device__ __forceinline__ hypot_sqr_func(const hypot_sqr_func& other) : binary_function<T, T, float>(){}\r
 420         __device__ __forceinline__ hypot_sqr_func() : binary_function<T, T, float>(){}\r
 421     };\r
 422 \r
 423     // Saturate Cast Functor\r
 424     template <typename T, typename D> struct saturate_cast_func : unary_function<T, D>\r
 425     {\r
 426         __device__ __forceinline__ D operator ()(typename TypeTraits<T>::ParameterType v) const\r
 427         {\r
 428             return saturate_cast<D>(v);\r
 429         }\r
 430         __device__ __forceinline__ saturate_cast_func(const saturate_cast_func& other):unary_function<T, D>(){}\r
 431         __device__ __forceinline__ saturate_cast_func():unary_function<T, D>(){}\r
 432     };\r
 433 \r
 434     // Threshold Functors\r
 435     template <typename T> struct thresh_binary_func : unary_function<T, T>\r
 436     {\r
 437         __host__ __device__ __forceinline__ thresh_binary_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {}\r
 438 \r
 439         __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const\r
 440         {\r
 441             return (src > thresh) * maxVal;\r
 442         }\r
 443 \r
 444         __device__ __forceinline__ thresh_binary_func(const thresh_binary_func& other)\r
 445             : unary_function<T, T>(), thresh(other.thresh), maxVal(other.maxVal){}\r
 446 \r
 447         __device__ __forceinline__ thresh_binary_func():unary_function<T, T>(){}\r
 448 \r
 449         const T thresh;\r
 450         const T maxVal;\r
 451     };\r
 452 \r
 453     template <typename T> struct thresh_binary_inv_func : unary_function<T, T>\r
 454     {\r
 455         __host__ __device__ __forceinline__ thresh_binary_inv_func(T thresh_, T maxVal_) : thresh(thresh_), maxVal(maxVal_) {}\r
 456 \r
 457         __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const\r
 458         {\r
 459             return (src <= thresh) * maxVal;\r
 460         }\r
 461 \r
 462         __device__ __forceinline__ thresh_binary_inv_func(const thresh_binary_inv_func& other)\r
 463             : unary_function<T, T>(), thresh(other.thresh), maxVal(other.maxVal){}\r
 464 \r
 465         __device__ __forceinline__ thresh_binary_inv_func():unary_function<T, T>(){}\r
 466 \r
 467         const T thresh;\r
 468         const T maxVal;\r
 469     };\r
 470 \r
 471     template <typename T> struct thresh_trunc_func : unary_function<T, T>\r
 472     {\r
 473         explicit __host__ __device__ __forceinline__ thresh_trunc_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;}\r
 474 \r
 475         __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const\r
 476         {\r
 477             return minimum<T>()(src, thresh);\r
 478         }\r
 479 \r
 480         __device__ __forceinline__ thresh_trunc_func(const thresh_trunc_func& other)\r
 481             : unary_function<T, T>(), thresh(other.thresh){}\r
 482 \r
 483         __device__ __forceinline__ thresh_trunc_func():unary_function<T, T>(){}\r
 484 \r
 485         const T thresh;\r
 486     };\r
 487 \r
 488     template <typename T> struct thresh_to_zero_func : unary_function<T, T>\r
 489     {\r
 490         explicit __host__ __device__ __forceinline__ thresh_to_zero_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;}\r
 491 \r
 492         __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const\r
 493         {\r
 494             return (src > thresh) * src;\r
 495         }\r
 496         __device__ __forceinline__ thresh_to_zero_func(const thresh_to_zero_func& other)\r
 497             : unary_function<T, T>(), thresh(other.thresh){}\r
 498 \r
 499         __device__ __forceinline__ thresh_to_zero_func():unary_function<T, T>(){}\r
 500 \r
 501         const T thresh;\r
 502     };\r
 503 \r
 504     template <typename T> struct thresh_to_zero_inv_func : unary_function<T, T>\r
 505     {\r
 506         explicit __host__ __device__ __forceinline__ thresh_to_zero_inv_func(T thresh_, T maxVal_ = 0) : thresh(thresh_) {(void)maxVal_;}\r
 507 \r
 508         __device__ __forceinline__ T operator()(typename TypeTraits<T>::ParameterType src) const\r
 509         {\r
 510             return (src <= thresh) * src;\r
 511         }\r
 512         __device__ __forceinline__ thresh_to_zero_inv_func(const thresh_to_zero_inv_func& other)\r
 513             : unary_function<T, T>(), thresh(other.thresh){}\r
 514 \r
 515         __device__ __forceinline__ thresh_to_zero_inv_func():unary_function<T, T>(){}\r
 516 \r
 517         const T thresh;\r
 518     };\r
 519 //bound!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!!! ============>\r
 520     // Function Object Adaptors\r
 521     template <typename Predicate> struct unary_negate : unary_function<typename Predicate::argument_type, bool>\r
 522     {\r
 523       explicit __host__ __device__ __forceinline__ unary_negate(const Predicate& p) : pred(p) {}\r
 524 \r
 525       __device__ __forceinline__ bool operator()(typename TypeTraits<typename Predicate::argument_type>::ParameterType x) const\r
 526       {\r
 527           return !pred(x);\r
 528       }\r
 529 \r
 530         __device__ __forceinline__ unary_negate(const unary_negate& other) : unary_function<typename Predicate::argument_type, bool>(){}\r
 531         __device__ __forceinline__ unary_negate() : unary_function<typename Predicate::argument_type, bool>(){}\r
 532 \r
 533       const Predicate pred;\r
 534     };\r
 535 \r
 536     template <typename Predicate> __host__ __device__ __forceinline__ unary_negate<Predicate> not1(const Predicate& pred)\r
 537     {\r
 538         return unary_negate<Predicate>(pred);\r
 539     }\r
 540 \r
 541     template <typename Predicate> struct binary_negate : binary_function<typename Predicate::first_argument_type, typename Predicate::second_argument_type, bool>\r
 542     {\r
 543         explicit __host__ __device__ __forceinline__ binary_negate(const Predicate& p) : pred(p) {}\r
 544 \r
 545         __device__ __forceinline__ bool operator()(typename TypeTraits<typename Predicate::first_argument_type>::ParameterType x,\r
 546                                                    typename TypeTraits<typename Predicate::second_argument_type>::ParameterType y) const\r
 547         {\r
 548             return !pred(x,y);\r
 549         }\r
 550         __device__ __forceinline__ binary_negate(const binary_negate& other)\r
 551         : binary_function<typename Predicate::first_argument_type, typename Predicate::second_argument_type, bool>(){}\r
 552 \r
 553         __device__ __forceinline__ binary_negate() :\r
 554         binary_function<typename Predicate::first_argument_type, typename Predicate::second_argument_type, bool>(){}\r
 555 \r
 556         const Predicate pred;\r
 557     };\r
 558 \r
 559     template <typename BinaryPredicate> __host__ __device__ __forceinline__ binary_negate<BinaryPredicate> not2(const BinaryPredicate& pred)\r
 560     {\r
 561         return binary_negate<BinaryPredicate>(pred);\r
 562     }\r
 563 \r
 564     template <typename Op> struct binder1st : unary_function<typename Op::second_argument_type, typename Op::result_type>\r
 565     {\r
 566         __host__ __device__ __forceinline__ binder1st(const Op& op_, const typename Op::first_argument_type& arg1_) : op(op_), arg1(arg1_) {}\r
 567 \r
 568         __device__ __forceinline__ typename Op::result_type operator ()(typename TypeTraits<typename Op::second_argument_type>::ParameterType a) const\r
 569         {\r
 570             return op(arg1, a);\r
 571         }\r
 572 \r
 573         __device__ __forceinline__ binder1st(const binder1st& other) :\r
 574         unary_function<typename Op::second_argument_type, typename Op::result_type>(){}\r
 575 \r
 576         const Op op;\r
 577         const typename Op::first_argument_type arg1;\r
 578     };\r
 579 \r
 580     template <typename Op, typename T> __host__ __device__ __forceinline__ binder1st<Op> bind1st(const Op& op, const T& x)\r
 581     {\r
 582         return binder1st<Op>(op, typename Op::first_argument_type(x));\r
 583     }\r
 584 \r
 585     template <typename Op> struct binder2nd : unary_function<typename Op::first_argument_type, typename Op::result_type>\r
 586     {\r
 587         __host__ __device__ __forceinline__ binder2nd(const Op& op_, const typename Op::second_argument_type& arg2_) : op(op_), arg2(arg2_) {}\r
 588 \r
 589         __forceinline__ __device__ typename Op::result_type operator ()(typename TypeTraits<typename Op::first_argument_type>::ParameterType a) const\r
 590         {\r
 591             return op(a, arg2);\r
 592         }\r
 593 \r
 594          __device__ __forceinline__ binder2nd(const binder2nd& other) :\r
 595         unary_function<typename Op::first_argument_type, typename Op::result_type>(), op(other.op), arg2(other.arg2){}\r
 596 \r
 597         const Op op;\r
 598         const typename Op::second_argument_type arg2;\r
 599     };\r
 600 \r
 601     template <typename Op, typename T> __host__ __device__ __forceinline__ binder2nd<Op> bind2nd(const Op& op, const T& x)\r
 602     {\r
 603         return binder2nd<Op>(op, typename Op::second_argument_type(x));\r
 604     }\r
 605 \r
 606     // Functor Traits\r
 607     template <typename F> struct IsUnaryFunction\r
 608     {\r
 609         typedef char Yes;\r
 610         struct No {Yes a[2];};\r
 611 \r
 612         template <typename T, typename D> static Yes check(unary_function<T, D>);\r
 613         static No check(...);\r
 614 \r
 615         static F makeF();\r
 616 \r
 617         enum { value = (sizeof(check(makeF())) == sizeof(Yes)) };\r
 618     };\r
 619 \r
 620     template <typename F> struct IsBinaryFunction\r
 621     {\r
 622         typedef char Yes;\r
 623         struct No {Yes a[2];};\r
 624 \r
 625         template <typename T1, typename T2, typename D> static Yes check(binary_function<T1, T2, D>);\r
 626         static No check(...);\r
 627 \r
 628         static F makeF();\r
 629 \r
 630         enum { value = (sizeof(check(makeF())) == sizeof(Yes)) };\r
 631     };\r
 632 \r
 633     namespace functional_detail\r
 634     {\r
 635         template <size_t src_elem_size, size_t dst_elem_size> struct UnOpShift { enum { shift = 1 }; };\r
 636         template <size_t src_elem_size> struct UnOpShift<src_elem_size, 1> { enum { shift = 4 }; };\r
 637         template <size_t src_elem_size> struct UnOpShift<src_elem_size, 2> { enum { shift = 2 }; };\r
 638 \r
 639         template <typename T, typename D> struct DefaultUnaryShift\r
 640         {\r
 641             enum { shift = UnOpShift<sizeof(T), sizeof(D)>::shift };\r
 642         };\r
 643 \r
 644         template <size_t src_elem_size1, size_t src_elem_size2, size_t dst_elem_size> struct BinOpShift { enum { shift = 1 }; };\r
 645         template <size_t src_elem_size1, size_t src_elem_size2> struct BinOpShift<src_elem_size1, src_elem_size2, 1> { enum { shift = 4 }; };\r
 646         template <size_t src_elem_size1, size_t src_elem_size2> struct BinOpShift<src_elem_size1, src_elem_size2, 2> { enum { shift = 2 }; };\r
 647 \r
 648         template <typename T1, typename T2, typename D> struct DefaultBinaryShift\r
 649         {\r
 650             enum { shift = BinOpShift<sizeof(T1), sizeof(T2), sizeof(D)>::shift };\r
 651         };\r
 652 \r
 653         template <typename Func, bool unary = IsUnaryFunction<Func>::value> struct ShiftDispatcher;\r
 654         template <typename Func> struct ShiftDispatcher<Func, true>\r
 655         {\r
 656             enum { shift = DefaultUnaryShift<typename Func::argument_type, typename Func::result_type>::shift };\r
 657         };\r
 658         template <typename Func> struct ShiftDispatcher<Func, false>\r
 659         {\r
 660             enum { shift = DefaultBinaryShift<typename Func::first_argument_type, typename Func::second_argument_type, typename Func::result_type>::shift };\r
 661         };\r
 662     }\r
 663 \r
 664     template <typename Func> struct DefaultTransformShift\r
 665     {\r
 666         enum { shift = functional_detail::ShiftDispatcher<Func>::shift };\r
 667     };\r
 668 \r
 669     template <typename Func> struct DefaultTransformFunctorTraits\r
 670     {\r
 671         enum { simple_block_dim_x = 16 };\r
 672         enum { simple_block_dim_y = 16 };\r
 673 \r
 674         enum { smart_block_dim_x = 16 };\r
 675         enum { smart_block_dim_y = 16 };\r
 676         enum { smart_shift = DefaultTransformShift<Func>::shift };\r
 677     };\r
 678 \r
 679     template <typename Func> struct TransformFunctorTraits : DefaultTransformFunctorTraits<Func> {};\r
 680 \r
 681 #define OPENCV_GPU_TRANSFORM_FUNCTOR_TRAITS(type) \\r
 682     template <> struct TransformFunctorTraits< type > : DefaultTransformFunctorTraits< type >\r
 683 }}} // namespace cv { namespace gpu { namespace device\r
 684 \r
 685 #endif // __OPENCV_GPU_FUNCTIONAL_HPP__\r