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