utility.hpp

00001 /*M///////////////////////////////////////////////////////////////////////////////////////
00002 //
00003 //  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
00004 //
00005 //  By downloading, copying, installing or using the software you agree to this license.
00006 //  If you do not agree to this license, do not download, install,
00007 //  copy or use the software.
00008 //
00009 //
00010 //                           License Agreement
00011 //                For Open Source Computer Vision Library
00012 //
00013 // Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
00014 // Copyright (C) 2009, Willow Garage Inc., all rights reserved.
00015 // Third party copyrights are property of their respective owners.
00016 //
00017 // Redistribution and use in source and binary forms, with or without modification,
00018 // are permitted provided that the following conditions are met:
00019 //
00020 //   * Redistribution's of source code must retain the above copyright notice,
00021 //     this list of conditions and the following disclaimer.
00022 //
00023 //   * Redistribution's in binary form must reproduce the above copyright notice,
00024 //     this list of conditions and the following disclaimer in the documentation
00025 //     and/or other materials provided with the distribution.
00026 //
00027 //   * The name of the copyright holders may not be used to endorse or promote products
00028 //     derived from this software without specific prior written permission.
00029 //
00030 // This software is provided by the copyright holders and contributors "as is" and
00031 // any express or implied warranties, including, but not limited to, the implied
00032 // warranties of merchantability and fitness for a particular purpose are disclaimed.
00033 // In no event shall the Intel Corporation or contributors be liable for any direct,
00034 // indirect, incidental, special, exemplary, or consequential damages
00035 // (including, but not limited to, procurement of substitute goods or services;
00036 // loss of use, data, or profits; or business interruption) however caused
00037 // and on any theory of liability, whether in contract, strict liability,
00038 // or tort (including negligence or otherwise) arising in any way out of
00039 // the use of this software, even if advised of the possibility of such damage.
00040 //
00041 //M*/
00042 
00043 #ifndef __OPENCV_CUDA_UTILITY_HPP__
00044 #define __OPENCV_CUDA_UTILITY_HPP__
00045 
00046 #include "saturate_cast.hpp "
00047 #include "datamov_utils.hpp "
00048 
00049 /** @file
00050  * @deprecated Use @ref cudev instead.
00051  */
00052 
00053 //! @cond IGNORED
00054 
00055 namespace cv { namespace cuda { namespace device
00056 {
00057     #define OPENCV_CUDA_LOG_WARP_SIZE        (5)
00058     #define OPENCV_CUDA_WARP_SIZE            (1 << OPENCV_CUDA_LOG_WARP_SIZE)
00059     #define OPENCV_CUDA_LOG_MEM_BANKS        ((__CUDA_ARCH__ >= 200) ? 5 : 4) // 32 banks on fermi, 16 on tesla
00060     #define OPENCV_CUDA_MEM_BANKS            (1 << OPENCV_CUDA_LOG_MEM_BANKS)
00061 
00062     ///////////////////////////////////////////////////////////////////////////////
00063     // swap
00064 
00065     template <typename T> void __device__ __host__ __forceinline__ swap(T& a, T& b)
00066     {
00067         const T temp = a;
00068         a = b;
00069         b = temp;
00070     }
00071 
00072     ///////////////////////////////////////////////////////////////////////////////
00073     // Mask Reader
00074 
00075     struct SingleMask
00076     {
00077         explicit __host__ __device__ __forceinline__ SingleMask(PtrStepb mask_) : mask(mask_) {}
00078         __host__ __device__ __forceinline__ SingleMask(const SingleMask& mask_): mask(mask_.mask){}
00079 
00080         __device__ __forceinline__ bool operator()(int y, int x) const
00081         {
00082             return mask.ptr(y)[x] != 0;
00083         }
00084 
00085         PtrStepb mask;
00086     };
00087 
00088     struct SingleMaskChannels
00089     {
00090         __host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_)
00091         : mask(mask_), channels(channels_) {}
00092         __host__ __device__ __forceinline__ SingleMaskChannels(const SingleMaskChannels& mask_)
00093             :mask(mask_.mask), channels(mask_.channels){}
00094 
00095         __device__ __forceinline__ bool operator()(int y, int x) const
00096         {
00097             return mask.ptr(y)[x / channels] != 0;
00098         }
00099 
00100         PtrStepb mask;
00101         int channels;
00102     };
00103 
00104     struct MaskCollection
00105     {
00106         explicit __host__ __device__ __forceinline__ MaskCollection(PtrStepb* maskCollection_)
00107             : maskCollection(maskCollection_) {}
00108 
00109         __device__ __forceinline__ MaskCollection(const MaskCollection& masks_)
00110             : maskCollection(masks_.maskCollection), curMask(masks_.curMask){}
00111 
00112         __device__ __forceinline__ void next()
00113         {
00114             curMask = *maskCollection++;
00115         }
00116         __device__ __forceinline__ void setMask(int z)
00117         {
00118             curMask = maskCollection[z];
00119         }
00120 
00121         __device__ __forceinline__ bool operator()(int y, int x) const
00122         {
00123             uchar val;
00124             return curMask.data == 0 || (ForceGlob<uchar>::Load(curMask.ptr(y), x, val), (val != 0));
00125         }
00126 
00127         const PtrStepb* maskCollection;
00128         PtrStepb curMask;
00129     };
00130 
00131     struct WithOutMask
00132     {
00133         __host__ __device__ __forceinline__ WithOutMask(){}
00134         __host__ __device__ __forceinline__ WithOutMask(const WithOutMask&){}
00135 
00136         __device__ __forceinline__ void next() const
00137         {
00138         }
00139         __device__ __forceinline__ void setMask(int) const
00140         {
00141         }
00142 
00143         __device__ __forceinline__ bool operator()(int, int) const
00144         {
00145             return true;
00146         }
00147 
00148         __device__ __forceinline__ bool operator()(int, int, int) const
00149         {
00150             return true;
00151         }
00152 
00153         static __device__ __forceinline__ bool check(int, int)
00154         {
00155             return true;
00156         }
00157 
00158         static __device__ __forceinline__ bool check(int, int, int)
00159         {
00160             return true;
00161         }
00162     };
00163 
00164     ///////////////////////////////////////////////////////////////////////////////
00165     // Solve linear system
00166 
00167     // solve 2x2 linear system Ax=b
00168     template <typename T> __device__ __forceinline__ bool solve2x2(const T A[2][2], const T b[2], T x[2])
00169     {
00170         T det = A[0][0] * A[1][1] - A[1][0] * A[0][1];
00171 
00172         if (det != 0)
00173         {
00174             double invdet = 1.0 / det;
00175 
00176             x[0] = saturate_cast<T>(invdet * (b[0] * A[1][1] - b[1] * A[0][1]));
00177 
00178             x[1] = saturate_cast<T>(invdet * (A[0][0] * b[1] - A[1][0] * b[0]));
00179 
00180             return true;
00181         }
00182 
00183         return false;
00184     }
00185 
00186     // solve 3x3 linear system Ax=b
00187     template <typename T> __device__ __forceinline__ bool solve3x3(const T A[3][3], const T b[3], T x[3])
00188     {
00189         T det = A[0][0] * (A[1][1] * A[2][2] - A[1][2] * A[2][1])
00190               - A[0][1] * (A[1][0] * A[2][2] - A[1][2] * A[2][0])
00191               + A[0][2] * (A[1][0] * A[2][1] - A[1][1] * A[2][0]);
00192 
00193         if (det != 0)
00194         {
00195             double invdet = 1.0 / det;
00196 
00197             x[0] = saturate_cast<T>(invdet *
00198                 (b[0]    * (A[1][1] * A[2][2] - A[1][2] * A[2][1]) -
00199                  A[0][1] * (b[1]    * A[2][2] - A[1][2] * b[2]   ) +
00200                  A[0][2] * (b[1]    * A[2][1] - A[1][1] * b[2]   )));
00201 
00202             x[1] = saturate_cast<T>(invdet *
00203                 (A[0][0] * (b[1]    * A[2][2] - A[1][2] * b[2]   ) -
00204                  b[0]    * (A[1][0] * A[2][2] - A[1][2] * A[2][0]) +
00205                  A[0][2] * (A[1][0] * b[2]    - b[1]    * A[2][0])));
00206 
00207             x[2] = saturate_cast<T>(invdet *
00208                 (A[0][0] * (A[1][1] * b[2]    - b[1]    * A[2][1]) -
00209                  A[0][1] * (A[1][0] * b[2]    - b[1]    * A[2][0]) +
00210                  b[0]    * (A[1][0] * A[2][1] - A[1][1] * A[2][0])));
00211 
00212             return true;
00213         }
00214 
00215         return false;
00216     }
00217 }}} // namespace cv { namespace cuda { namespace cudev
00218 
00219 //! @endcond
00220 
00221 #endif // __OPENCV_CUDA_UTILITY_HPP__
00222