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