utility.hpp
Go to the documentation of this file.
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_GPU_UTILITY_HPP__
44 #define __OPENCV_GPU_UTILITY_HPP__
45 
46 #include "saturate_cast.hpp"
47 #include "datamov_utils.hpp"
48 
49 namespace cv { namespace gpu { namespace device
50 {
51  #define OPENCV_GPU_LOG_WARP_SIZE (5)
52  #define OPENCV_GPU_WARP_SIZE (1 << OPENCV_GPU_LOG_WARP_SIZE)
53  #define OPENCV_GPU_LOG_MEM_BANKS ((__CUDA_ARCH__ >= 200) ? 5 : 4) // 32 banks on fermi, 16 on tesla
54  #define OPENCV_GPU_MEM_BANKS (1 << OPENCV_GPU_LOG_MEM_BANKS)
55 
57  // swap
58 
59  template <typename T> void __device__ __host__ __forceinline__ swap(T& a, T& b)
60  {
61  const T temp = a;
62  a = b;
63  b = temp;
64  }
65 
67  // Mask Reader
68 
69  struct SingleMask
70  {
71  explicit __host__ __device__ __forceinline__ SingleMask(PtrStepb mask_) : mask(mask_) {}
72  __host__ __device__ __forceinline__ SingleMask(const SingleMask& mask_): mask(mask_.mask){}
73 
74  __device__ __forceinline__ bool operator()(int y, int x) const
75  {
76  return mask.ptr(y)[x] != 0;
77  }
78 
80  };
81 
83  {
84  __host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_)
85  : mask(mask_), channels(channels_) {}
86  __host__ __device__ __forceinline__ SingleMaskChannels(const SingleMaskChannels& mask_)
87  :mask(mask_.mask), channels(mask_.channels){}
88 
89  __device__ __forceinline__ bool operator()(int y, int x) const
90  {
91  return mask.ptr(y)[x / channels] != 0;
92  }
93 
95  int channels;
96  };
97 
99  {
100  explicit __host__ __device__ __forceinline__ MaskCollection(PtrStepb* maskCollection_)
101  : maskCollection(maskCollection_) {}
102 
103  __device__ __forceinline__ MaskCollection(const MaskCollection& masks_)
104  : maskCollection(masks_.maskCollection), curMask(masks_.curMask){}
105 
106  __device__ __forceinline__ void next()
107  {
108  curMask = *maskCollection++;
109  }
110  __device__ __forceinline__ void setMask(int z)
111  {
112  curMask = maskCollection[z];
113  }
114 
115  __device__ __forceinline__ bool operator()(int y, int x) const
116  {
117  uchar val;
118  return curMask.data == 0 || (ForceGlob<uchar>::Load(curMask.ptr(y), x, val), (val != 0));
119  }
120 
123  };
124 
125  struct WithOutMask
126  {
127  __host__ __device__ __forceinline__ WithOutMask(){}
128  __host__ __device__ __forceinline__ WithOutMask(const WithOutMask&){}
129 
130  __device__ __forceinline__ void next() const
131  {
132  }
133  __device__ __forceinline__ void setMask(int) const
134  {
135  }
136 
137  __device__ __forceinline__ bool operator()(int, int) const
138  {
139  return true;
140  }
141 
142  __device__ __forceinline__ bool operator()(int, int, int) const
143  {
144  return true;
145  }
146 
147  static __device__ __forceinline__ bool check(int, int)
148  {
149  return true;
150  }
151 
152  static __device__ __forceinline__ bool check(int, int, int)
153  {
154  return true;
155  }
156  };
157 
159  // Solve linear system
160 
161  // solve 2x2 linear system Ax=b
162  template <typename T> __device__ __forceinline__ bool solve2x2(const T A[2][2], const T b[2], T x[2])
163  {
164  T det = A[0][0] * A[1][1] - A[1][0] * A[0][1];
165 
166  if (det != 0)
167  {
168  double invdet = 1.0 / det;
169 
170  x[0] = saturate_cast<T>(invdet * (b[0] * A[1][1] - b[1] * A[0][1]));
171 
172  x[1] = saturate_cast<T>(invdet * (A[0][0] * b[1] - A[1][0] * b[0]));
173 
174  return true;
175  }
176 
177  return false;
178  }
179 
180  // solve 3x3 linear system Ax=b
181  template <typename T> __device__ __forceinline__ bool solve3x3(const T A[3][3], const T b[3], T x[3])
182  {
183  T det = A[0][0] * (A[1][1] * A[2][2] - A[1][2] * A[2][1])
184  - A[0][1] * (A[1][0] * A[2][2] - A[1][2] * A[2][0])
185  + A[0][2] * (A[1][0] * A[2][1] - A[1][1] * A[2][0]);
186 
187  if (det != 0)
188  {
189  double invdet = 1.0 / det;
190 
191  x[0] = saturate_cast<T>(invdet *
192  (b[0] * (A[1][1] * A[2][2] - A[1][2] * A[2][1]) -
193  A[0][1] * (b[1] * A[2][2] - A[1][2] * b[2] ) +
194  A[0][2] * (b[1] * A[2][1] - A[1][1] * b[2] )));
195 
196  x[1] = saturate_cast<T>(invdet *
197  (A[0][0] * (b[1] * A[2][2] - A[1][2] * b[2] ) -
198  b[0] * (A[1][0] * A[2][2] - A[1][2] * A[2][0]) +
199  A[0][2] * (A[1][0] * b[2] - b[1] * A[2][0])));
200 
201  x[2] = saturate_cast<T>(invdet *
202  (A[0][0] * (A[1][1] * b[2] - b[1] * A[2][1]) -
203  A[0][1] * (A[1][0] * b[2] - b[1] * A[2][0]) +
204  b[0] * (A[1][0] * A[2][1] - A[1][1] * A[2][0])));
205 
206  return true;
207  }
208 
209  return false;
210  }
211 }}} // namespace cv { namespace gpu { namespace device
212 
213 #endif // __OPENCV_GPU_UTILITY_HPP__
GLenum GLint GLint y
Definition: core_c.h:613
__host__ __device__ __forceinline__ WithOutMask()
Definition: utility.hpp:127
static __device__ __forceinline__ bool check(int, int, int)
Definition: utility.hpp:152
__device__ __forceinline__ MaskCollection(const MaskCollection &masks_)
Definition: utility.hpp:103
__device__ __forceinline__ void setMask(int) const
Definition: utility.hpp:133
PtrStepb curMask
Definition: utility.hpp:122
__host__ __device__ __forceinline__ MaskCollection(PtrStepb *maskCollection_)
Definition: utility.hpp:100
Definition: utility.hpp:125
PtrStepb mask
Definition: utility.hpp:79
__device__ __forceinline__ void next() const
Definition: utility.hpp:130
CvPoint2D32f float float b
Definition: legacy.hpp:578
const CvMat const CvMat const CvMat CvMat CvMat CvMat CvMat CvSize CvMat CvMat * T
Definition: calib3d.hpp:270
__host__ __device__ __forceinline__ SingleMask(PtrStepb mask_)
Definition: utility.hpp:71
__device__ __forceinline__ bool operator()(int, int) const
Definition: utility.hpp:137
__host__ __device__ __forceinline__ SingleMask(const SingleMask &mask_)
Definition: utility.hpp:72
static __device__ __forceinline__ bool check(int, int)
Definition: utility.hpp:147
unsigned char uchar
Definition: common.hpp:100
GLuint GLfloat * val
__host__ __device__ __forceinline__ SingleMaskChannels(const SingleMaskChannels &mask_)
Definition: utility.hpp:86
__device__ __forceinline__ bool solve3x3(const T A[3][3], const T b[3], T x[3])
Definition: utility.hpp:181
__device__ __forceinline__ void next()
Definition: utility.hpp:106
__device__ __forceinline__ bool solve2x2(const T A[2][2], const T b[2], T x[2])
Definition: utility.hpp:162
GLenum GLint x
Definition: core_c.h:632
void __device__ __host__ __forceinline__ swap(T &a, T &b)
Definition: utility.hpp:59
GLdouble GLdouble z
GLboolean GLboolean GLboolean b
Definition: legacy.hpp:633
__device__ __forceinline__ _Tp saturate_cast(uchar v)
Definition: saturate_cast.hpp:50
__host__ __device__ __forceinline__ WithOutMask(const WithOutMask &)
Definition: utility.hpp:128
__device__ __forceinline__ bool operator()(int, int, int) const
Definition: utility.hpp:142
__device__ __forceinline__ bool operator()(int y, int x) const
Definition: utility.hpp:89
Definition: utility.hpp:98
PtrStep< unsigned char > PtrStepb
Definition: cuda_devptrs.hpp:121
int channels
Definition: utility.hpp:95
GLboolean GLboolean GLboolean GLboolean a
Definition: legacy.hpp:633
__device__ __forceinline__ void setMask(int z)
Definition: utility.hpp:110
const PtrStepb * maskCollection
Definition: utility.hpp:121
CvPoint2D32f float a
Definition: legacy.hpp:578
int x
Definition: highgui_c.h:186
PtrStepb mask
Definition: utility.hpp:94
CvArr CvArr * temp
Definition: imgproc_c.h:242
__device__ __forceinline__ bool operator()(int y, int x) const
Definition: utility.hpp:74
GLenum GLint GLuint mask
Definition: tracking.hpp:132
__device__ __forceinline__ bool operator()(int y, int x) const
Definition: utility.hpp:115
__host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_)
Definition: utility.hpp:84
__device__ static __forceinline__ void Load(const T *ptr, int offset, T &val)
Definition: datamov_utils.hpp:55
Definition: utility.hpp:82
Definition: utility.hpp:69