43 #ifndef __OPENCV_GPU_UTILITY_HPP__
44 #define __OPENCV_GPU_UTILITY_HPP__
49 namespace cv {
namespace gpu {
namespace device
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)
59 template <
typename T>
void __device__ __host__ __forceinline__
swap(
T&
a,
T&
b)
76 return mask.ptr(y)[
x] != 0;
106 __device__ __forceinline__
void next()
130 __device__ __forceinline__
void next()
const
133 __device__ __forceinline__
void setMask(
int)
const
142 __device__ __forceinline__
bool operator()(
int,
int,
int)
const
147 static __device__ __forceinline__
bool check(
int,
int)
152 static __device__ __forceinline__
bool check(
int,
int,
int)
162 template <
typename T> __device__ __forceinline__
bool solve2x2(
const T A[2][2],
const T b[2],
T x[2])
164 T det = A[0][0] * A[1][1] - A[1][0] * A[0][1];
168 double invdet = 1.0 / det;
170 x[0] =
saturate_cast<
T>(invdet * (b[0] * A[1][1] - b[1] * A[0][1]));
172 x[1] =
saturate_cast<T>(invdet * (A[0][0] * b[1] - A[1][0] * b[0]));
181 template <
typename T> __device__ __forceinline__
bool solve3x3(
const T A[3][3],
const T b[3],
T x[3])
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]);
189 double invdet = 1.0 / det;
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] )));
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])));
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])));
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
__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
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