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_UTILITY_HPP__ 44 #define __OPENCV_CUDA_UTILITY_HPP__ 45 46 #include "saturate_cast.hpp" 47 #include "datamov_utils.hpp" 48 49 /** @file 50 * @deprecated Use @ref cudev instead. 51 */ 52 53 //! @cond IGNORED 54 55 namespace cv { namespace cuda { namespace device 56 { 57 #define OPENCV_CUDA_LOG_WARP_SIZE (5) 58 #define OPENCV_CUDA_WARP_SIZE (1 << OPENCV_CUDA_LOG_WARP_SIZE) 59 #define OPENCV_CUDA_LOG_MEM_BANKS ((__CUDA_ARCH__ >= 200) ? 5 : 4) // 32 banks on fermi, 16 on tesla 60 #define OPENCV_CUDA_MEM_BANKS (1 << OPENCV_CUDA_LOG_MEM_BANKS) 61 62 /////////////////////////////////////////////////////////////////////////////// 63 // swap 64 swap(T & a,T & b)65 template <typename T> void __device__ __host__ __forceinline__ swap(T& a, T& b) 66 { 67 const T temp = a; 68 a = b; 69 b = temp; 70 } 71 72 /////////////////////////////////////////////////////////////////////////////// 73 // Mask Reader 74 75 struct SingleMask 76 { SingleMaskcv::cuda::device::SingleMask77 explicit __host__ __device__ __forceinline__ SingleMask(PtrStepb mask_) : mask(mask_) {} SingleMaskcv::cuda::device::SingleMask78 __host__ __device__ __forceinline__ SingleMask(const SingleMask& mask_): mask(mask_.mask){} 79 operator ()cv::cuda::device::SingleMask80 __device__ __forceinline__ bool operator()(int y, int x) const 81 { 82 return mask.ptr(y)[x] != 0; 83 } 84 85 PtrStepb mask; 86 }; 87 88 struct SingleMaskChannels 89 { SingleMaskChannelscv::cuda::device::SingleMaskChannels90 __host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_) 91 : mask(mask_), channels(channels_) {} SingleMaskChannelscv::cuda::device::SingleMaskChannels92 __host__ __device__ __forceinline__ SingleMaskChannels(const SingleMaskChannels& mask_) 93 :mask(mask_.mask), channels(mask_.channels){} 94 operator ()cv::cuda::device::SingleMaskChannels95 __device__ __forceinline__ bool operator()(int y, int x) const 96 { 97 return mask.ptr(y)[x / channels] != 0; 98 } 99 100 PtrStepb mask; 101 int channels; 102 }; 103 104 struct MaskCollection 105 { MaskCollectioncv::cuda::device::MaskCollection106 explicit __host__ __device__ __forceinline__ MaskCollection(PtrStepb* maskCollection_) 107 : maskCollection(maskCollection_) {} 108 MaskCollectioncv::cuda::device::MaskCollection109 __device__ __forceinline__ MaskCollection(const MaskCollection& masks_) 110 : maskCollection(masks_.maskCollection), curMask(masks_.curMask){} 111 nextcv::cuda::device::MaskCollection112 __device__ __forceinline__ void next() 113 { 114 curMask = *maskCollection++; 115 } setMaskcv::cuda::device::MaskCollection116 __device__ __forceinline__ void setMask(int z) 117 { 118 curMask = maskCollection[z]; 119 } 120 operator ()cv::cuda::device::MaskCollection121 __device__ __forceinline__ bool operator()(int y, int x) const 122 { 123 uchar val; 124 return curMask.data == 0 || (ForceGlob<uchar>::Load(curMask.ptr(y), x, val), (val != 0)); 125 } 126 127 const PtrStepb* maskCollection; 128 PtrStepb curMask; 129 }; 130 131 struct WithOutMask 132 { WithOutMaskcv::cuda::device::WithOutMask133 __host__ __device__ __forceinline__ WithOutMask(){} WithOutMaskcv::cuda::device::WithOutMask134 __host__ __device__ __forceinline__ WithOutMask(const WithOutMask&){} 135 nextcv::cuda::device::WithOutMask136 __device__ __forceinline__ void next() const 137 { 138 } setMaskcv::cuda::device::WithOutMask139 __device__ __forceinline__ void setMask(int) const 140 { 141 } 142 operator ()cv::cuda::device::WithOutMask143 __device__ __forceinline__ bool operator()(int, int) const 144 { 145 return true; 146 } 147 operator ()cv::cuda::device::WithOutMask148 __device__ __forceinline__ bool operator()(int, int, int) const 149 { 150 return true; 151 } 152 checkcv::cuda::device::WithOutMask153 static __device__ __forceinline__ bool check(int, int) 154 { 155 return true; 156 } 157 checkcv::cuda::device::WithOutMask158 static __device__ __forceinline__ bool check(int, int, int) 159 { 160 return true; 161 } 162 }; 163 164 /////////////////////////////////////////////////////////////////////////////// 165 // Solve linear system 166 167 // solve 2x2 linear system Ax=b solve2x2(const T A[2][2],const T b[2],T x[2])168 template <typename T> __device__ __forceinline__ bool solve2x2(const T A[2][2], const T b[2], T x[2]) 169 { 170 T det = A[0][0] * A[1][1] - A[1][0] * A[0][1]; 171 172 if (det != 0) 173 { 174 double invdet = 1.0 / det; 175 176 x[0] = saturate_cast<T>(invdet * (b[0] * A[1][1] - b[1] * A[0][1])); 177 178 x[1] = saturate_cast<T>(invdet * (A[0][0] * b[1] - A[1][0] * b[0])); 179 180 return true; 181 } 182 183 return false; 184 } 185 186 // solve 3x3 linear system Ax=b solve3x3(const T A[3][3],const T b[3],T x[3])187 template <typename T> __device__ __forceinline__ bool solve3x3(const T A[3][3], const T b[3], T x[3]) 188 { 189 T det = A[0][0] * (A[1][1] * A[2][2] - A[1][2] * A[2][1]) 190 - A[0][1] * (A[1][0] * A[2][2] - A[1][2] * A[2][0]) 191 + A[0][2] * (A[1][0] * A[2][1] - A[1][1] * A[2][0]); 192 193 if (det != 0) 194 { 195 double invdet = 1.0 / det; 196 197 x[0] = saturate_cast<T>(invdet * 198 (b[0] * (A[1][1] * A[2][2] - A[1][2] * A[2][1]) - 199 A[0][1] * (b[1] * A[2][2] - A[1][2] * b[2] ) + 200 A[0][2] * (b[1] * A[2][1] - A[1][1] * b[2] ))); 201 202 x[1] = saturate_cast<T>(invdet * 203 (A[0][0] * (b[1] * A[2][2] - A[1][2] * b[2] ) - 204 b[0] * (A[1][0] * A[2][2] - A[1][2] * A[2][0]) + 205 A[0][2] * (A[1][0] * b[2] - b[1] * A[2][0]))); 206 207 x[2] = saturate_cast<T>(invdet * 208 (A[0][0] * (A[1][1] * b[2] - b[1] * A[2][1]) - 209 A[0][1] * (A[1][0] * b[2] - b[1] * A[2][0]) + 210 b[0] * (A[1][0] * A[2][1] - A[1][1] * A[2][0]))); 211 212 return true; 213 } 214 215 return false; 216 } 217 }}} // namespace cv { namespace cuda { namespace cudev 218 219 //! @endcond 220 221 #endif // __OPENCV_CUDA_UTILITY_HPP__ 222