123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230 |
- #ifndef OPENCV_CUDA_UTILITY_HPP
- #define OPENCV_CUDA_UTILITY_HPP
- #include "saturate_cast.hpp"
- #include "datamov_utils.hpp"
- namespace cv { namespace cuda { namespace device
- {
- struct CV_EXPORTS ThrustAllocator
- {
- typedef uchar value_type;
- virtual ~ThrustAllocator();
- virtual __device__ __host__ uchar* allocate(size_t numBytes) = 0;
- virtual __device__ __host__ void deallocate(uchar* ptr, size_t numBytes) = 0;
- static ThrustAllocator& getAllocator();
- static void setAllocator(ThrustAllocator* allocator);
- };
- #define OPENCV_CUDA_LOG_WARP_SIZE (5)
- #define OPENCV_CUDA_WARP_SIZE (1 << OPENCV_CUDA_LOG_WARP_SIZE)
- #define OPENCV_CUDA_LOG_MEM_BANKS ((__CUDA_ARCH__ >= 200) ? 5 : 4)
- #define OPENCV_CUDA_MEM_BANKS (1 << OPENCV_CUDA_LOG_MEM_BANKS)
-
-
- template <typename T> void __device__ __host__ __forceinline__ swap(T& a, T& b)
- {
- const T temp = a;
- a = b;
- b = temp;
- }
-
-
- struct SingleMask
- {
- explicit __host__ __device__ __forceinline__ SingleMask(PtrStepb mask_) : mask(mask_) {}
- __host__ __device__ __forceinline__ SingleMask(const SingleMask& mask_): mask(mask_.mask){}
- __device__ __forceinline__ bool operator()(int y, int x) const
- {
- return mask.ptr(y)[x] != 0;
- }
- PtrStepb mask;
- };
- struct SingleMaskChannels
- {
- __host__ __device__ __forceinline__ SingleMaskChannels(PtrStepb mask_, int channels_)
- : mask(mask_), channels(channels_) {}
- __host__ __device__ __forceinline__ SingleMaskChannels(const SingleMaskChannels& mask_)
- :mask(mask_.mask), channels(mask_.channels){}
- __device__ __forceinline__ bool operator()(int y, int x) const
- {
- return mask.ptr(y)[x / channels] != 0;
- }
- PtrStepb mask;
- int channels;
- };
- struct MaskCollection
- {
- explicit __host__ __device__ __forceinline__ MaskCollection(PtrStepb* maskCollection_)
- : maskCollection(maskCollection_) {}
- __device__ __forceinline__ MaskCollection(const MaskCollection& masks_)
- : maskCollection(masks_.maskCollection), curMask(masks_.curMask){}
- __device__ __forceinline__ void next()
- {
- curMask = *maskCollection++;
- }
- __device__ __forceinline__ void setMask(int z)
- {
- curMask = maskCollection[z];
- }
- __device__ __forceinline__ bool operator()(int y, int x) const
- {
- uchar val;
- return curMask.data == 0 || (ForceGlob<uchar>::Load(curMask.ptr(y), x, val), (val != 0));
- }
- const PtrStepb* maskCollection;
- PtrStepb curMask;
- };
- struct WithOutMask
- {
- __host__ __device__ __forceinline__ WithOutMask(){}
- __host__ __device__ __forceinline__ WithOutMask(const WithOutMask&){}
- __device__ __forceinline__ void next() const
- {
- }
- __device__ __forceinline__ void setMask(int) const
- {
- }
- __device__ __forceinline__ bool operator()(int, int) const
- {
- return true;
- }
- __device__ __forceinline__ bool operator()(int, int, int) const
- {
- return true;
- }
- static __device__ __forceinline__ bool check(int, int)
- {
- return true;
- }
- static __device__ __forceinline__ bool check(int, int, int)
- {
- return true;
- }
- };
-
-
-
- template <typename T> __device__ __forceinline__ bool solve2x2(const T A[2][2], const T b[2], T x[2])
- {
- T det = A[0][0] * A[1][1] - A[1][0] * A[0][1];
- if (det != 0)
- {
- double invdet = 1.0 / det;
- x[0] = saturate_cast<T>(invdet * (b[0] * A[1][1] - b[1] * A[0][1]));
- x[1] = saturate_cast<T>(invdet * (A[0][0] * b[1] - A[1][0] * b[0]));
- return true;
- }
- return false;
- }
-
- template <typename T> __device__ __forceinline__ bool solve3x3(const T A[3][3], const T b[3], T x[3])
- {
- T det = A[0][0] * (A[1][1] * A[2][2] - A[1][2] * A[2][1])
- - A[0][1] * (A[1][0] * A[2][2] - A[1][2] * A[2][0])
- + A[0][2] * (A[1][0] * A[2][1] - A[1][1] * A[2][0]);
- if (det != 0)
- {
- double invdet = 1.0 / det;
- x[0] = saturate_cast<T>(invdet *
- (b[0] * (A[1][1] * A[2][2] - A[1][2] * A[2][1]) -
- A[0][1] * (b[1] * A[2][2] - A[1][2] * b[2] ) +
- A[0][2] * (b[1] * A[2][1] - A[1][1] * b[2] )));
- x[1] = saturate_cast<T>(invdet *
- (A[0][0] * (b[1] * A[2][2] - A[1][2] * b[2] ) -
- b[0] * (A[1][0] * A[2][2] - A[1][2] * A[2][0]) +
- A[0][2] * (A[1][0] * b[2] - b[1] * A[2][0])));
- x[2] = saturate_cast<T>(invdet *
- (A[0][0] * (A[1][1] * b[2] - b[1] * A[2][1]) -
- A[0][1] * (A[1][0] * b[2] - b[1] * A[2][0]) +
- b[0] * (A[1][0] * A[2][1] - A[1][1] * A[2][0])));
- return true;
- }
- return false;
- }
- }}}
- #endif
|