Home | History | Annotate | Download | only in cuda
      1 /*M///////////////////////////////////////////////////////////////////////////////////////
      2 //
      3 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
      4 //
      5 //  By downloading, copying, installing or using the software you agree to this license.
      6 //  If you do not agree to this license, do not download, install,
      7 //  copy or use the software.
      8 //
      9 //
     10 //                           License Agreement
     11 //                For Open Source Computer Vision Library
     12 //
     13 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
     14 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
     15 // Third party copyrights are property of their respective owners.
     16 //
     17 // Redistribution and use in source and binary forms, with or without modification,
     18 // are permitted provided that the following conditions are met:
     19 //
     20 //   * Redistribution's of source code must retain the above copyright notice,
     21 //     this list of conditions and the following disclaimer.
     22 //
     23 //   * Redistribution's in binary form must reproduce the above copyright notice,
     24 //     this list of conditions and the following disclaimer in the documentation
     25 //     and/or other materials provided with the distribution.
     26 //
     27 //   * The name of the copyright holders may not be used to endorse or promote products
     28 //     derived from this software without specific prior written permission.
     29 //
     30 // This software is provided by the copyright holders and contributors "as is" and
     31 // any express or implied warranties, including, but not limited to, the implied
     32 // warranties of merchantability and fitness for a particular purpose are disclaimed.
     33 // In no event shall the Intel Corporation or contributors be liable for any direct,
     34 // indirect, incidental, special, exemplary, or consequential damages
     35 // (including, but not limited to, procurement of substitute goods or services;
     36 // loss of use, data, or profits; or business interruption) however caused
     37 // and on any theory of liability, whether in contract, strict liability,
     38 // or tort (including negligence or otherwise) arising in any way out of
     39 // the use of this software, even if advised of the possibility of such damage.
     40 //
     41 //M*/
     42 
     43 #ifndef __OPENCV_CUDA_FILTERS_HPP__
     44 #define __OPENCV_CUDA_FILTERS_HPP__
     45 
     46 #include "saturate_cast.hpp"
     47 #include "vec_traits.hpp"
     48 #include "vec_math.hpp"
     49 #include "type_traits.hpp"
     50 
     51 /** @file
     52  * @deprecated Use @ref cudev instead.
     53  */
     54 
     55 //! @cond IGNORED
     56 
     57 namespace cv { namespace cuda { namespace device
     58 {
     59     template <typename Ptr2D> struct PointFilter
     60     {
     61         typedef typename Ptr2D::elem_type elem_type;
     62         typedef float index_type;
     63 
     64         explicit __host__ __device__ __forceinline__ PointFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f)
     65         : src(src_)
     66         {
     67             (void)fx;
     68             (void)fy;
     69         }
     70 
     71         __device__ __forceinline__ elem_type operator ()(float y, float x) const
     72         {
     73             return src(__float2int_rz(y), __float2int_rz(x));
     74         }
     75 
     76         Ptr2D src;
     77     };
     78 
     79     template <typename Ptr2D> struct LinearFilter
     80     {
     81         typedef typename Ptr2D::elem_type elem_type;
     82         typedef float index_type;
     83 
     84         explicit __host__ __device__ __forceinline__ LinearFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f)
     85         : src(src_)
     86         {
     87             (void)fx;
     88             (void)fy;
     89         }
     90         __device__ __forceinline__ elem_type operator ()(float y, float x) const
     91         {
     92             typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type;
     93 
     94             work_type out = VecTraits<work_type>::all(0);
     95 
     96             const int x1 = __float2int_rd(x);
     97             const int y1 = __float2int_rd(y);
     98             const int x2 = x1 + 1;
     99             const int y2 = y1 + 1;
    100 
    101             elem_type src_reg = src(y1, x1);
    102             out = out + src_reg * ((x2 - x) * (y2 - y));
    103 
    104             src_reg = src(y1, x2);
    105             out = out + src_reg * ((x - x1) * (y2 - y));
    106 
    107             src_reg = src(y2, x1);
    108             out = out + src_reg * ((x2 - x) * (y - y1));
    109 
    110             src_reg = src(y2, x2);
    111             out = out + src_reg * ((x - x1) * (y - y1));
    112 
    113             return saturate_cast<elem_type>(out);
    114         }
    115 
    116         Ptr2D src;
    117     };
    118 
    119     template <typename Ptr2D> struct CubicFilter
    120     {
    121         typedef typename Ptr2D::elem_type elem_type;
    122         typedef float index_type;
    123         typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type;
    124 
    125         explicit __host__ __device__ __forceinline__ CubicFilter(const Ptr2D& src_, float fx = 0.f, float fy = 0.f)
    126         : src(src_)
    127         {
    128             (void)fx;
    129             (void)fy;
    130         }
    131 
    132         static __device__ __forceinline__ float bicubicCoeff(float x_)
    133         {
    134             float x = fabsf(x_);
    135             if (x <= 1.0f)
    136             {
    137                 return x * x * (1.5f * x - 2.5f) + 1.0f;
    138             }
    139             else if (x < 2.0f)
    140             {
    141                 return x * (x * (-0.5f * x + 2.5f) - 4.0f) + 2.0f;
    142             }
    143             else
    144             {
    145                 return 0.0f;
    146             }
    147         }
    148 
    149         __device__ elem_type operator ()(float y, float x) const
    150         {
    151             const float xmin = ::ceilf(x - 2.0f);
    152             const float xmax = ::floorf(x + 2.0f);
    153 
    154             const float ymin = ::ceilf(y - 2.0f);
    155             const float ymax = ::floorf(y + 2.0f);
    156 
    157             work_type sum = VecTraits<work_type>::all(0);
    158             float wsum = 0.0f;
    159 
    160             for (float cy = ymin; cy <= ymax; cy += 1.0f)
    161             {
    162                 for (float cx = xmin; cx <= xmax; cx += 1.0f)
    163                 {
    164                     const float w = bicubicCoeff(x - cx) * bicubicCoeff(y - cy);
    165                     sum = sum + w * src(__float2int_rd(cy), __float2int_rd(cx));
    166                     wsum += w;
    167                 }
    168             }
    169 
    170             work_type res = (!wsum)? VecTraits<work_type>::all(0) : sum / wsum;
    171 
    172             return saturate_cast<elem_type>(res);
    173         }
    174 
    175         Ptr2D src;
    176     };
    177     // for integer scaling
    178     template <typename Ptr2D> struct IntegerAreaFilter
    179     {
    180         typedef typename Ptr2D::elem_type elem_type;
    181         typedef float index_type;
    182 
    183         explicit __host__ __device__ __forceinline__ IntegerAreaFilter(const Ptr2D& src_, float scale_x_, float scale_y_)
    184             : src(src_), scale_x(scale_x_), scale_y(scale_y_), scale(1.f / (scale_x * scale_y)) {}
    185 
    186         __device__ __forceinline__ elem_type operator ()(float y, float x) const
    187         {
    188             float fsx1 = x * scale_x;
    189             float fsx2 = fsx1 + scale_x;
    190 
    191             int sx1 = __float2int_ru(fsx1);
    192             int sx2 = __float2int_rd(fsx2);
    193 
    194             float fsy1 = y * scale_y;
    195             float fsy2 = fsy1 + scale_y;
    196 
    197             int sy1 = __float2int_ru(fsy1);
    198             int sy2 = __float2int_rd(fsy2);
    199 
    200             typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type;
    201             work_type out = VecTraits<work_type>::all(0.f);
    202 
    203             for(int dy = sy1; dy < sy2; ++dy)
    204                 for(int dx = sx1; dx < sx2; ++dx)
    205                 {
    206                     out = out + src(dy, dx) * scale;
    207                 }
    208 
    209             return saturate_cast<elem_type>(out);
    210         }
    211 
    212         Ptr2D src;
    213         float scale_x, scale_y ,scale;
    214     };
    215 
    216     template <typename Ptr2D> struct AreaFilter
    217     {
    218         typedef typename Ptr2D::elem_type elem_type;
    219         typedef float index_type;
    220 
    221         explicit __host__ __device__ __forceinline__ AreaFilter(const Ptr2D& src_, float scale_x_, float scale_y_)
    222             : src(src_), scale_x(scale_x_), scale_y(scale_y_){}
    223 
    224         __device__ __forceinline__ elem_type operator ()(float y, float x) const
    225         {
    226             float fsx1 = x * scale_x;
    227             float fsx2 = fsx1 + scale_x;
    228 
    229             int sx1 = __float2int_ru(fsx1);
    230             int sx2 = __float2int_rd(fsx2);
    231 
    232             float fsy1 = y * scale_y;
    233             float fsy2 = fsy1 + scale_y;
    234 
    235             int sy1 = __float2int_ru(fsy1);
    236             int sy2 = __float2int_rd(fsy2);
    237 
    238             float scale = 1.f / (fminf(scale_x, src.width - fsx1) * fminf(scale_y, src.height - fsy1));
    239 
    240             typedef typename TypeVec<float, VecTraits<elem_type>::cn>::vec_type work_type;
    241             work_type out = VecTraits<work_type>::all(0.f);
    242 
    243             for (int dy = sy1; dy < sy2; ++dy)
    244             {
    245                 for (int dx = sx1; dx < sx2; ++dx)
    246                     out = out + src(dy, dx) * scale;
    247 
    248                 if (sx1 > fsx1)
    249                     out = out + src(dy, (sx1 -1) ) * ((sx1 - fsx1) * scale);
    250 
    251                 if (sx2 < fsx2)
    252                     out = out + src(dy, sx2) * ((fsx2 -sx2) * scale);
    253             }
    254 
    255             if (sy1 > fsy1)
    256                 for (int dx = sx1; dx < sx2; ++dx)
    257                     out = out + src( (sy1 - 1) , dx) * ((sy1 -fsy1) * scale);
    258 
    259             if (sy2 < fsy2)
    260                 for (int dx = sx1; dx < sx2; ++dx)
    261                     out = out + src(sy2, dx) * ((fsy2 -sy2) * scale);
    262 
    263             if ((sy1 > fsy1) &&  (sx1 > fsx1))
    264                 out = out + src( (sy1 - 1) , (sx1 - 1)) * ((sy1 -fsy1) * (sx1 -fsx1) * scale);
    265 
    266             if ((sy1 > fsy1) &&  (sx2 < fsx2))
    267                 out = out + src( (sy1 - 1) , sx2) * ((sy1 -fsy1) * (fsx2 -sx2) * scale);
    268 
    269             if ((sy2 < fsy2) &&  (sx2 < fsx2))
    270                 out = out + src(sy2, sx2) * ((fsy2 -sy2) * (fsx2 -sx2) * scale);
    271 
    272             if ((sy2 < fsy2) &&  (sx1 > fsx1))
    273                 out = out + src(sy2, (sx1 - 1)) * ((fsy2 -sy2) * (sx1 -fsx1) * scale);
    274 
    275             return saturate_cast<elem_type>(out);
    276         }
    277 
    278         Ptr2D src;
    279         float scale_x, scale_y;
    280         int width, haight;
    281     };
    282 }}} // namespace cv { namespace cuda { namespace cudev
    283 
    284 //! @endcond
    285 
    286 #endif // __OPENCV_CUDA_FILTERS_HPP__
    287