sz.height = src1.rows;\r
\r
int funcIdx = normType >> 1;\r
- Scalar retVal;\r
+ double retVal;\r
\r
nppSafeCall( npp_norm_diff_func[funcIdx](src1.ptr<Npp8u>(), src1.step, \r
src2.ptr<Npp8u>(), src2.step, \r
- sz, retVal.val) );\r
+ sz, &retVal) );\r
\r
- return retVal[0];\r
+ return retVal;\r
}\r
\r
////////////////////////////////////////////////////////////////////////\r
\r
Scalar cv::gpu::sum(const GpuMat& src)\r
{\r
- CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);\r
- \r
- \r
- \r
+ CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4); \r
\r
NppiSize sz;\r
sz.width = src.cols;\r
GpuMat buf(1, bufsz, CV_32S);\r
\r
Scalar res;\r
- nppSafeCall( nppiSum_8u_C1R(src.ptr<Npp8u>(), src.step, sz, buf.ptr<Npp32s>(), res.val) );\r
+ nppSafeCall( nppiSum_8u_C1R(src.ptr<Npp8u>(), src.step, sz, buf.ptr<Npp32s>(), res.val) );\r
return res;\r
}\r
else\r
nppSafeCall( nppiSum_8u_C4R(src.ptr<Npp8u>(), src.step, sz, buf.ptr<Npp32s>(), res.val) );\r
return res;\r
}\r
-\r
- \r
}\r
\r
////////////////////////////////////////////////////////////////////////\r
{\r
public:\r
Npp32s pLevels[256];\r
+ const Npp32s* pLevels3[3];\r
+ int nValues3[3];\r
\r
LevelsInit()\r
- { \r
+ {\r
+ nValues3[0] = nValues3[1] = nValues3[2] = 256;\r
for (int i = 0; i < 256; ++i)\r
pLevels[i] = i;\r
+ pLevels3[0] = pLevels3[1] = pLevels3[2] = pLevels;\r
}\r
};\r
static LevelsInit lvls;\r
\r
int cn = src.channels();\r
\r
- CV_Assert(src.type() == CV_8UC1);\r
- CV_Assert(lut.depth() == CV_32SC1 && lut.rows * lut.cols == 256 && lut.isContinuous());\r
+ CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC3);\r
+ CV_Assert(lut.depth() == CV_8U && (lut.channels() == 1 || lut.channels() == cn) && lut.rows * lut.cols == 256 && lut.isContinuous());\r
\r
- dst.create(src.size(), src.type());\r
+ dst.create(src.size(), CV_MAKETYPE(lut.depth(), cn));\r
\r
NppiSize sz;\r
sz.height = src.rows;\r
sz.width = src.cols;\r
+ \r
+ Mat nppLut;\r
+ lut.convertTo(nppLut, CV_32S);\r
\r
- nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr<Npp8u>(), src.step, dst.ptr<Npp8u>(), dst.step, sz, \r
- lut.ptr<Npp32s>(), lvls.pLevels, 256) );\r
+ if (src.type() == CV_8UC1)\r
+ {\r
+ nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr<Npp8u>(), src.step, dst.ptr<Npp8u>(), dst.step, sz, \r
+ nppLut.ptr<Npp32s>(), lvls.pLevels, 256) );\r
+ }\r
+ else\r
+ {\r
+ Mat nppLut3[3];\r
+ const Npp32s* pValues3[3];\r
+ if (nppLut.channels() == 1)\r
+ pValues3[0] = pValues3[1] = pValues3[2] = nppLut.ptr<Npp32s>();\r
+ else\r
+ {\r
+ cv::split(nppLut, nppLut3);\r
+ pValues3[0] = nppLut3[0].ptr<Npp32s>();\r
+ pValues3[1] = nppLut3[1].ptr<Npp32s>(); \r
+ pValues3[2] = nppLut3[2].ptr<Npp32s>();\r
+ }\r
+ nppSafeCall( nppiLUT_Linear_8u_C3R(src.ptr<Npp8u>(), src.step, dst.ptr<Npp8u>(), dst.step, sz, \r
+ pValues3, lvls.pLevels3, lvls.nValues3) );\r
+ }\r
}\r
\r
#endif /* !defined (HAVE_CUDA) */
\ No newline at end of file
template<> struct TypeVec<float, 3> { typedef float3 vec_t; };\r
template<> struct TypeVec<float, 4> { typedef float4 vec_t; };\r
\r
- template<typename _Tp> struct ColorChannel {};\r
+ template<typename T> struct ColorChannel {};\r
\r
template<> struct ColorChannel<uchar>\r
{\r
typedef float worktype_f;\r
static __device__ float max() { return 1.f; }\r
static __device__ float half() { return 0.5f; }\r
- }; \r
+ };\r
+\r
+ template <typename T>\r
+ __device__ void assignAlpha(typename TypeVec<T, 3>::vec_t& vec, T val)\r
+ {\r
+ }\r
+ template <typename T>\r
+ __device__ void assignAlpha(typename TypeVec<T, 4>::vec_t& vec, T val)\r
+ {\r
+ vec.w = val;\r
+ }\r
}\r
\r
//////////////////////////////////////// SwapChannels /////////////////////////////////////\r
__constant__ int ccoeffs[4];\r
\r
template <int CN, typename T>\r
- __global__ void swapChannels(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols)\r
+ __global__ void swapChannels(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols)\r
{\r
typedef typename TypeVec<T, CN>::vec_t vec_t;\r
\r
\r
namespace cv { namespace gpu { namespace improc\r
{\r
- template <typename T>\r
- void swapChannels_caller(const DevMem2D_<T>& src, const DevMem2D_<T>& dst, int cn, const int* coeffs, cudaStream_t stream)\r
+ template <typename T, int CN>\r
+ void swapChannels_caller(const DevMem2D& src, const DevMem2D& dst, const int* coeffs, cudaStream_t stream)\r
{\r
dim3 threads(32, 8, 1);\r
dim3 grid(1, 1, 1);\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, cn * sizeof(int)) );\r
+ cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, CN * sizeof(int)) );\r
\r
- switch (cn)\r
- {\r
- case 3:\r
- imgproc::swapChannels<3><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
- break;\r
- case 4:\r
- imgproc::swapChannels<4><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
- break;\r
- default:\r
- cv::gpu::error("Unsupported channels count", __FILE__, __LINE__);\r
- break;\r
- }\r
+ imgproc::swapChannels<CN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ dst.ptr, dst.step, src.rows, src.cols);\r
\r
if (stream == 0)\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
\r
- void swapChannels_gpu(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream)\r
+ void swapChannels_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream)\r
{\r
- swapChannels_caller(src, dst, cn, coeffs, stream);\r
+ typedef void (*swapChannels_caller_t)(const DevMem2D& src, const DevMem2D& dst, const int* coeffs, cudaStream_t stream);\r
+ static const swapChannels_caller_t swapChannels_callers[] = {swapChannels_caller<uchar, 3>, swapChannels_caller<uchar, 4>};\r
+\r
+ swapChannels_callers[cn - 3](src, dst, coeffs, stream);\r
}\r
\r
- void swapChannels_gpu(const DevMem2D_<unsigned short>& src, const DevMem2D_<unsigned short>& dst, int cn, const int* coeffs, cudaStream_t stream)\r
+ void swapChannels_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream)\r
{\r
- swapChannels_caller(src, dst, cn, coeffs, stream);\r
+ typedef void (*swapChannels_caller_t)(const DevMem2D& src, const DevMem2D& dst, const int* coeffs, cudaStream_t stream);\r
+ static const swapChannels_caller_t swapChannels_callers[] = {swapChannels_caller<unsigned short, 3>, swapChannels_caller<unsigned short, 4>};\r
+\r
+ swapChannels_callers[cn - 3](src, dst, coeffs, stream);\r
}\r
\r
- void swapChannels_gpu(const DevMem2Df& src, const DevMem2Df& dst, int cn, const int* coeffs, cudaStream_t stream)\r
+ void swapChannels_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream)\r
{\r
- swapChannels_caller(src, dst, cn, coeffs, stream);\r
- }\r
+ typedef void (*swapChannels_caller_t)(const DevMem2D& src, const DevMem2D& dst, const int* coeffs, cudaStream_t stream);\r
+ static const swapChannels_caller_t swapChannels_callers[] = {swapChannels_caller<float, 3>, swapChannels_caller<float, 4>};\r
+\r
+ swapChannels_callers[cn - 3](src, dst, coeffs, stream);\r
+ } \r
}}}\r
\r
////////////////// Various 3/4-channel to 3/4-channel RGB transformations /////////////////\r
namespace imgproc\r
{\r
template <int SRCCN, int DSTCN, typename T>\r
- __global__ void RGB2RGB(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx)\r
+ __global__ void RGB2RGB(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx)\r
{\r
typedef typename TypeVec<T, SRCCN>::vec_t src_t;\r
typedef typename TypeVec<T, DSTCN>::vec_t dst_t;\r
dst.x = ((const T*)(&src))[bidx];\r
dst.y = src.y;\r
dst.z = ((const T*)(&src))[bidx ^ 2];\r
- if (DSTCN == 4)\r
- ((T*)(&dst))[3] = ColorChannel<T>::max();\r
+ assignAlpha(dst, ColorChannel<T>::max());\r
\r
*(dst_t*)(dst_ + y * dst_step + x * DSTCN) = dst;\r
}\r
\r
namespace cv { namespace gpu { namespace improc\r
{\r
- template <typename T>\r
- void RGB2RGB_caller(const DevMem2D_<T>& src, int srccn, const DevMem2D_<T>& dst, int dstcn, int bidx, cudaStream_t stream)\r
+ template <typename T, int SRCCN, int DSTCN>\r
+ void RGB2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)\r
{\r
dim3 threads(32, 8, 1);\r
dim3 grid(1, 1, 1);\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- switch (dstcn)\r
- {\r
- case 3:\r
- switch (srccn)\r
- {\r
- case 3:\r
- {\r
- int coeffs[] = {2, 1, 0};\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, 3 * sizeof(int)) );\r
- imgproc::swapChannels<3><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
- break;\r
- }\r
- case 4:\r
- imgproc::RGB2RGB<4, 3><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T),\r
- src.rows, src.cols, bidx);\r
- break;\r
- default:\r
- cv::gpu::error("Unsupported channels count", __FILE__, __LINE__);\r
- break;\r
- }\r
- break;\r
- case 4:\r
- switch (srccn)\r
- {\r
- case 3:\r
- imgproc::RGB2RGB<3, 4><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T),\r
- src.rows, src.cols, bidx);\r
- break;\r
- case 4:\r
- {\r
- int coeffs[] = {2, 1, 0, 3};\r
- cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, 4 * sizeof(int)) );\r
- imgproc::swapChannels<4><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
- break;\r
- }\r
- default:\r
- cv::gpu::error("Unsupported channels count", __FILE__, __LINE__);\r
- break;\r
- }\r
- break;\r
- default:\r
- cv::gpu::error("Unsupported channels count", __FILE__, __LINE__);\r
- break;\r
- }\r
+ imgproc::RGB2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ dst.ptr, dst.step, src.rows, src.cols, bidx);\r
\r
if (stream == 0)\r
cudaSafeCall( cudaThreadSynchronize() );\r
}\r
\r
- void RGB2RGB_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream)\r
+ void RGB2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream)\r
{\r
- RGB2RGB_caller(src, srccn, dst, dstcn, bidx, stream);\r
+ typedef void (*RGB2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
+ static const RGB2RGB_caller_t RGB2RGB_callers[2][2] = \r
+ {\r
+ {RGB2RGB_caller<uchar, 3, 3>, RGB2RGB_caller<uchar, 3, 4>}, \r
+ {RGB2RGB_caller<uchar, 4, 3>, RGB2RGB_caller<uchar, 4, 4>}\r
+ };\r
+\r
+ RGB2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
}\r
\r
- void RGB2RGB_gpu(const DevMem2D_<unsigned short>& src, int srccn, const DevMem2D_<unsigned short>& dst, int dstcn, int bidx, cudaStream_t stream)\r
+ void RGB2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream)\r
{\r
- RGB2RGB_caller(src, srccn, dst, dstcn, bidx, stream);\r
+ typedef void (*RGB2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
+ static const RGB2RGB_caller_t RGB2RGB_callers[2][2] = \r
+ {\r
+ {RGB2RGB_caller<unsigned short, 3, 3>, RGB2RGB_caller<unsigned short, 3, 4>}, \r
+ {RGB2RGB_caller<unsigned short, 4, 3>, RGB2RGB_caller<unsigned short, 4, 4>}\r
+ };\r
+\r
+ RGB2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
}\r
\r
- void RGB2RGB_gpu(const DevMem2Df& src, int srccn, const DevMem2Df& dst, int dstcn, int bidx, cudaStream_t stream)\r
+ void RGB2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream)\r
{\r
- RGB2RGB_caller(src, srccn, dst, dstcn, bidx, stream);\r
+ typedef void (*RGB2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
+ static const RGB2RGB_caller_t RGB2RGB_callers[2][2] = \r
+ {\r
+ {RGB2RGB_caller<float, 3, 3>, RGB2RGB_caller<float, 3, 4>}, \r
+ {RGB2RGB_caller<float, 4, 3>, RGB2RGB_caller<float, 4, 4>}\r
+ };\r
+\r
+ RGB2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
}\r
}}}\r
\r
/////////// Transforming 16-bit (565 or 555) RGB to/from 24/32-bit (888[8]) RGB //////////\r
\r
-//namespace imgproc\r
-//{\r
-// struct RGB5x52RGB\r
-// {\r
-// typedef uchar channel_type;\r
-//\r
-// RGB5x52RGB(int _dstcn, int _blueIdx, int _greenBits)\r
-// : dstcn(_dstcn), blueIdx(_blueIdx), greenBits(_greenBits) {}\r
-//\r
-// void operator()(const uchar* src, uchar* dst, int n) const\r
-// {\r
-// int dcn = dstcn, bidx = blueIdx;\r
-// if( greenBits == 6 )\r
-// for( int i = 0; i < n; i++, dst += dcn )\r
-// {\r
-// unsigned t = ((const unsigned short*)src)[i];\r
-// dst[bidx] = (uchar)(t << 3);\r
-// dst[1] = (uchar)((t >> 3) & ~3);\r
-// dst[bidx ^ 2] = (uchar)((t >> 8) & ~7);\r
-// if( dcn == 4 )\r
-// dst[3] = 255;\r
-// }\r
-// else\r
-// for( int i = 0; i < n; i++, dst += dcn )\r
-// {\r
-// unsigned t = ((const unsigned short*)src)[i];\r
-// dst[bidx] = (uchar)(t << 3);\r
-// dst[1] = (uchar)((t >> 2) & ~7);\r
-// dst[bidx ^ 2] = (uchar)((t >> 7) & ~7);\r
-// if( dcn == 4 )\r
-// dst[3] = t & 0x8000 ? 255 : 0;\r
-// }\r
-// }\r
-//\r
-// int dstcn, blueIdx, greenBits;\r
-// };\r
-//\r
-//\r
-// struct RGB2RGB5x5\r
-// {\r
-// typedef uchar channel_type;\r
-//\r
-// RGB2RGB5x5(int _srccn, int _blueIdx, int _greenBits)\r
-// : srccn(_srccn), blueIdx(_blueIdx), greenBits(_greenBits) {}\r
-//\r
-// void operator()(const uchar* src, uchar* dst, int n) const\r
-// {\r
-// int scn = srccn, bidx = blueIdx;\r
-// if( greenBits == 6 )\r
-// for( int i = 0; i < n; i++, src += scn )\r
-// {\r
-// ((unsigned short*)dst)[i] = (unsigned short)((src[bidx] >> 3)|((src[1]&~3) << 3)|((src[bidx^2]&~7) << 8));\r
-// }\r
-// else if( scn == 3 )\r
-// for( int i = 0; i < n; i++, src += 3 )\r
-// {\r
-// ((unsigned short*)dst)[i] = (unsigned short)((src[bidx] >> 3)|((src[1]&~7) << 2)|((src[bidx^2]&~7) << 7));\r
-// }\r
-// else\r
-// for( int i = 0; i < n; i++, src += 4 )\r
-// {\r
-// ((unsigned short*)dst)[i] = (unsigned short)((src[bidx] >> 3)|((src[1]&~7) << 2)|\r
-// ((src[bidx^2]&~7) << 7)|(src[3] ? 0x8000 : 0));\r
-// }\r
-// }\r
-//\r
-// int srccn, blueIdx, greenBits;\r
-// };\r
-//}\r
-//\r
-//namespace cv { namespace gpu { namespace impl\r
-//{\r
-//}}}\r
-\r
-///////////////////////////////// Grayscale to Color ////////////////////////////////\r
-\r
namespace imgproc\r
{\r
- template <typename T>\r
- __global__ void Gray2RGB_3(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols)\r
+ template <int GREEN_BITS, int DSTCN> struct RGB5x52RGBConverter {};\r
+ \r
+ template <int DSTCN> struct RGB5x52RGBConverter<5, DSTCN>\r
{\r
+ typedef typename TypeVec<uchar, DSTCN>::vec_t dst_t;\r
+\r
+ static __device__ dst_t cvt(unsigned int src, int bidx)\r
+ {\r
+ dst_t dst;\r
+ \r
+ ((uchar*)(&dst))[bidx] = (uchar)(src << 3);\r
+ dst.y = (uchar)((src >> 2) & ~7);\r
+ ((uchar*)(&dst))[bidx ^ 2] = (uchar)((src >> 7) & ~7);\r
+ assignAlpha(dst, (uchar)(src & 0x8000 ? 255 : 0));\r
+\r
+ return dst;\r
+ }\r
+ };\r
+ template <int DSTCN> struct RGB5x52RGBConverter<6, DSTCN>\r
+ {\r
+ typedef typename TypeVec<uchar, DSTCN>::vec_t dst_t;\r
+\r
+ static __device__ dst_t cvt(unsigned int src, int bidx)\r
+ {\r
+ dst_t dst;\r
+ \r
+ ((uchar*)(&dst))[bidx] = (uchar)(src << 3);\r
+ dst.y = (uchar)((src >> 3) & ~3);\r
+ ((uchar*)(&dst))[bidx ^ 2] = (uchar)((src >> 8) & ~7);\r
+ assignAlpha(dst, (uchar)(255));\r
+\r
+ return dst;\r
+ }\r
+ };\r
+\r
+ template <int GREEN_BITS, int DSTCN>\r
+ __global__ void RGB5x52RGB(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx)\r
+ {\r
+ typedef typename TypeVec<uchar, DSTCN>::vec_t dst_t;\r
+\r
const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
\r
if (y < rows && x < cols)\r
{\r
- T src = src_[y * src_step + x];\r
- T* dst = dst_ + y * dst_step + x * 3;\r
- dst[0] = src;\r
- dst[1] = src;\r
- dst[2] = src;\r
+ unsigned int src = *(const unsigned short*)(src_ + y * src_step + (x << 1));\r
+ \r
+ *(dst_t*)(dst_ + y * dst_step + x * DSTCN) = RGB5x52RGBConverter<GREEN_BITS, DSTCN>::cvt(src, bidx);\r
}\r
}\r
\r
- template <typename T>\r
- __global__ void Gray2RGB_4(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols)\r
+ /*struct RGB5x52RGB\r
+ {\r
+ typedef uchar channel_type;\r
+\r
+ RGB5x52RGB(int _dstcn, int _blueIdx, int _greenBits)\r
+ : dstcn(_dstcn), blueIdx(_blueIdx), greenBits(_greenBits) {}\r
+\r
+ void operator()(const uchar* src, uchar* dst, int n) const\r
+ {\r
+ int dcn = dstcn, bidx = blueIdx;\r
+ if( greenBits == 6 )\r
+ for( int i = 0; i < n; i++, dst += dcn )\r
+ {\r
+ unsigned t = ((const unsigned short*)src)[i];\r
+ dst[bidx] = (uchar)(t << 3);\r
+ dst[1] = (uchar)((t >> 3) & ~3);\r
+ dst[bidx ^ 2] = (uchar)((t >> 8) & ~7);\r
+ if( dcn == 4 )\r
+ dst[3] = 255;\r
+ }\r
+ else\r
+ for( int i = 0; i < n; i++, dst += dcn )\r
+ {\r
+ unsigned t = ((const unsigned short*)src)[i];\r
+ dst[bidx] = (uchar)(t << 3);\r
+ dst[1] = (uchar)((t >> 2) & ~7);\r
+ dst[bidx ^ 2] = (uchar)((t >> 7) & ~7);\r
+ if( dcn == 4 )\r
+ dst[3] = t & 0x8000 ? 255 : 0;\r
+ }\r
+ }\r
+\r
+ int dstcn, blueIdx, greenBits;\r
+ };*/\r
+\r
+ template <int SRCCN, int GREEN_BITS> struct RGB2RGB5x5Converter {};\r
+\r
+ template<int SRCCN> struct RGB2RGB5x5Converter<SRCCN, 6> \r
+ {\r
+ static __device__ unsigned short cvt(const uchar* src_ptr, int bidx)\r
+ {\r
+ return (unsigned short)((src_ptr[bidx] >> 3) | ((src_ptr[1] & ~3) << 3) | ((src_ptr[bidx^2] & ~7) << 8));\r
+ }\r
+ };\r
+ template<> struct RGB2RGB5x5Converter<3, 5> \r
+ {\r
+ static __device__ unsigned short cvt(const uchar* src_ptr, int bidx)\r
+ {\r
+ return (unsigned short)((src_ptr[bidx] >> 3) | ((src_ptr[1] & ~7) << 2) | ((src_ptr[bidx^2] & ~7) << 7));\r
+ }\r
+ };\r
+ template<> struct RGB2RGB5x5Converter<4, 5> \r
{\r
- typedef typename TypeVec<T, 4>::vec_t vec4_t;\r
+ static __device__ unsigned short cvt(const uchar* src_ptr, int bidx)\r
+ {\r
+ return (unsigned short)((src_ptr[bidx] >> 3) | ((src_ptr[1] & ~7) << 2) | ((src_ptr[bidx^2] & ~7) << 7)|(src_ptr[3] ? 0x8000 : 0));\r
+ }\r
+ }; \r
+\r
+ template<int SRCCN, int GREEN_BITS>\r
+ __global__ void RGB2RGB5x5(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx)\r
+ {\r
+ typedef typename TypeVec<uchar, SRCCN>::vec_t src_t;\r
+\r
+ const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
+ const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+\r
+ if (y < rows && x < cols)\r
+ {\r
+ src_t src = *(src_t*)(src_ + y * src_step + x * SRCCN);\r
+\r
+ *(unsigned short*)(dst_ + y * dst_step + (x << 1)) = RGB2RGB5x5Converter<SRCCN, GREEN_BITS>::cvt((const uchar*)(&src), bidx);\r
+ }\r
+ }\r
+}\r
+\r
+namespace cv { namespace gpu { namespace improc\r
+{\r
+ template <int GREEN_BITS, int DSTCN>\r
+ void RGB5x52RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)\r
+ {\r
+ dim3 threads(32, 8, 1);\r
+ dim3 grid(1, 1, 1);\r
+\r
+ grid.x = divUp(src.cols, threads.x);\r
+ grid.y = divUp(src.rows, threads.y);\r
+\r
+ imgproc::RGB5x52RGB<GREEN_BITS, DSTCN><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ dst.ptr, dst.step, src.rows, src.cols, bidx);\r
+\r
+ if (stream == 0)\r
+ cudaSafeCall( cudaThreadSynchronize() );\r
+ }\r
+\r
+ void RGB5x52RGB_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream)\r
+ {\r
+ typedef void (*RGB5x52RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
+ static const RGB5x52RGB_caller_t RGB5x52RGB_callers[2][2] = \r
+ {\r
+ {RGB5x52RGB_caller<5, 3>, RGB5x52RGB_caller<5, 4>},\r
+ {RGB5x52RGB_caller<6, 3>, RGB5x52RGB_caller<6, 4>}\r
+ };\r
+\r
+ RGB5x52RGB_callers[green_bits - 5][dstcn - 5](src, dst, bidx, stream);\r
+ }\r
+\r
+ template <int SRCCN, int GREEN_BITS>\r
+ void RGB2RGB5x5_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)\r
+ {\r
+ dim3 threads(32, 8, 1);\r
+ dim3 grid(1, 1, 1);\r
+\r
+ grid.x = divUp(src.cols, threads.x);\r
+ grid.y = divUp(src.rows, threads.y);\r
+\r
+ imgproc::RGB2RGB5x5<SRCCN, GREEN_BITS><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+ dst.ptr, dst.step, src.rows, src.cols, bidx);\r
+\r
+ if (stream == 0)\r
+ cudaSafeCall( cudaThreadSynchronize() );\r
+ }\r
+\r
+ void RGB2RGB5x5_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int green_bits, int bidx, cudaStream_t stream)\r
+ {\r
+ typedef void (*RGB2RGB5x5_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
+ static const RGB2RGB5x5_caller_t RGB2RGB5x5_callers[2][2] = \r
+ {\r
+ {RGB2RGB5x5_caller<3, 5>, RGB2RGB5x5_caller<3, 6>},\r
+ {RGB2RGB5x5_caller<4, 5>, RGB2RGB5x5_caller<4, 6>}\r
+ };\r
+\r
+ RGB2RGB5x5_callers[srccn - 3][green_bits - 5](src, dst, bidx, stream);\r
+ }\r
+}}}\r
+\r
+///////////////////////////////// Grayscale to Color ////////////////////////////////\r
+\r
+namespace imgproc\r
+{\r
+ template <int DSTCN, typename T>\r
+ __global__ void Gray2RGB(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols)\r
+ {\r
+ typedef typename TypeVec<T, DSTCN>::vec_t dst_t;\r
\r
const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
if (y < rows && x < cols)\r
{\r
T src = src_[y * src_step + x];\r
- vec4_t dst;\r
+ dst_t dst;\r
dst.x = src;\r
dst.y = src;\r
dst.z = src;\r
- dst.w = ColorChannel<T>::max();\r
- *(vec4_t*)(dst_ + y * dst_step + (x << 2)) = dst;\r
+ assignAlpha(dst, ColorChannel<T>::max());\r
+ *(dst_t*)(dst_ + y * dst_step + x * DSTCN) = dst;\r
}\r
}\r
\r
\r
namespace cv { namespace gpu { namespace improc\r
{\r
- template <typename T>\r
- void Gray2RGB_caller(const DevMem2D_<T>& src, const DevMem2D_<T>& dst, int dstcn, cudaStream_t stream)\r
+ template <typename T, int DSTCN>\r
+ void Gray2RGB_caller(const DevMem2D_<T>& src, const DevMem2D_<T>& dst, cudaStream_t stream)\r
{\r
dim3 threads(32, 8, 1);\r
dim3 grid(1, 1, 1);\r
grid.x = divUp(src.cols, threads.x);\r
grid.y = divUp(src.rows, threads.y);\r
\r
- switch (dstcn)\r
- {\r
- case 3:\r
- imgproc::Gray2RGB_3<<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
- break;\r
- case 4:\r
- imgproc::Gray2RGB_4<<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
- break;\r
- default:\r
- cv::gpu::error("Unsupported channels count", __FILE__, __LINE__);\r
- break;\r
- }\r
+ imgproc::Gray2RGB<DSTCN><<<grid, threads, 0, stream>>>(src.ptr, src.step / sizeof(T), \r
+ dst.ptr, dst.step / sizeof(T), src.rows, src.cols);\r
\r
if (stream == 0)\r
cudaSafeCall( cudaThreadSynchronize() );\r
\r
void Gray2RGB_gpu(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream)\r
{\r
- Gray2RGB_caller(src, dst, dstcn, stream);\r
+ typedef void (*Gray2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream);\r
+ static const Gray2RGB_caller_t Gray2RGB_callers[] = {Gray2RGB_caller<uchar, 3>, Gray2RGB_caller<uchar, 4>};\r
+\r
+ Gray2RGB_callers[dstcn - 3](src, dst, stream);\r
}\r
\r
void Gray2RGB_gpu(const DevMem2D_<unsigned short>& src, const DevMem2D_<unsigned short>& dst, int dstcn, cudaStream_t stream)\r
{\r
- Gray2RGB_caller(src, dst, dstcn, stream);\r
+ typedef void (*Gray2RGB_caller_t)(const DevMem2D_<unsigned short>& src, const DevMem2D_<unsigned short>& dst, cudaStream_t stream);\r
+ static const Gray2RGB_caller_t Gray2RGB_callers[] = {Gray2RGB_caller<unsigned short, 3>, Gray2RGB_caller<unsigned short, 4>};\r
+\r
+ Gray2RGB_callers[dstcn - 3](src, dst, stream);\r
}\r
\r
void Gray2RGB_gpu(const DevMem2Df& src, const DevMem2Df& dst, int dstcn, cudaStream_t stream)\r
{\r
- Gray2RGB_caller(src, dst, dstcn, stream);\r
+ typedef void (*Gray2RGB_caller_t)(const DevMem2Df& src, const DevMem2Df& dst, cudaStream_t stream);\r
+ static const Gray2RGB_caller_t Gray2RGB_callers[] = {Gray2RGB_caller<float, 3>, Gray2RGB_caller<float, 4>};\r
+\r
+ Gray2RGB_callers[dstcn - 3](src, dst, stream);\r
}\r
}}}\r
\r
void reprojectImageTo3D_gpu(const DevMem2D& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);\r
void reprojectImageTo3D_gpu(const DevMem2D_<short>& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream);\r
\r
- void swapChannels_gpu(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream);\r
- void swapChannels_gpu(const DevMem2D_<ushort>& src, const DevMem2D_<ushort>& dst, int cn, const int* coeffs, cudaStream_t stream);\r
- void swapChannels_gpu(const DevMem2Df& src, const DevMem2Df& dst, int cn, const int* coeffs, cudaStream_t stream);\r
+ void swapChannels_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream);\r
+ void swapChannels_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream);\r
+ void swapChannels_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream);\r
\r
- void RGB2RGB_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
- void RGB2RGB_gpu(const DevMem2D_<ushort>& src, int srccn, const DevMem2D_<ushort>& dst, int dstcn, int bidx, cudaStream_t stream);\r
- void RGB2RGB_gpu(const DevMem2Df& src, int srccn, const DevMem2Df& dst, int dstcn, int bidx, cudaStream_t stream);\r
+ void RGB2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
+ void RGB2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
+ void RGB2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
+\r
+ void RGB5x52RGB_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream);\r
+ void RGB2RGB5x5_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int green_bits, int bidx, cudaStream_t stream);\r
\r
void Gray2RGB_gpu(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream);\r
void Gray2RGB_gpu(const DevMem2D_<ushort>& src, const DevMem2D_<ushort>& dst, int dstcn, cudaStream_t stream);\r
\r
out.create(sz, CV_MAKETYPE(depth, dcn));\r
if( depth == CV_8U )\r
- improc::RGB2RGB_gpu((DevMem2D)src, scn, (DevMem2D)out, dcn, bidx, stream);\r
+ improc::RGB2RGB_gpu_8u(src, scn, out, dcn, bidx, stream);\r
else if( depth == CV_16U )\r
- improc::RGB2RGB_gpu((DevMem2D_<unsigned short>)src, scn, (DevMem2D_<unsigned short>)out, dcn, bidx, stream);\r
+ improc::RGB2RGB_gpu_16u(src, scn, out, dcn, bidx, stream);\r
else\r
- improc::RGB2RGB_gpu((DevMem2Df)src, scn, (DevMem2Df)out, dcn, bidx, stream);\r
+ improc::RGB2RGB_gpu_32f(src, scn, out, dcn, bidx, stream);\r
break;\r
\r
- //case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:\r
- //case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555:\r
- // CV_Assert( (scn == 3 || scn == 4) && depth == CV_8U );\r
- // dst.create(sz, CV_8UC2);\r
- //\r
- // CvtColorLoop(src, dst, RGB2RGB5x5(scn,\r
- // code == CV_BGR2BGR565 || code == CV_BGR2BGR555 ||\r
- // code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2,\r
- // code == CV_BGR2BGR565 || code == CV_RGB2BGR565 ||\r
- // code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5 // green bits\r
- // ));\r
- // break;\r
+ case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:\r
+ case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555:\r
+ CV_Assert( (scn == 3 || scn == 4) && depth == CV_8U );\r
+ out.create(sz, CV_8UC2);\r
+\r
+ improc::RGB2RGB5x5_gpu(src, scn, out, code == CV_BGR2BGR565 || code == CV_RGB2BGR565 ||\r
+ code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5,\r
+ code == CV_BGR2BGR565 || code == CV_BGR2BGR555 ||\r
+ code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2,\r
+ stream);\r
+ break;\r
\r
//case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:\r
//case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:\r
// if(dcn <= 0) dcn = 3;\r
// CV_Assert( (dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U );\r
- // dst.create(sz, CV_MAKETYPE(depth, dcn));\r
- // \r
- // CvtColorLoop(src, dst, RGB5x52RGB(dcn,\r
- // code == CV_BGR5652BGR || code == CV_BGR5552BGR ||\r
- // code == CV_BGR5652BGRA || code == CV_BGR5552BGRA ? 0 : 2, // blue idx\r
- // code == CV_BGR5652BGR || code == CV_BGR5652RGB ||\r
- // code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5 // green bits\r
- // ));\r
+ // out.create(sz, CV_MAKETYPE(depth, dcn));\r
+\r
+ // improc::RGB5x52RGB_gpu(src, code == CV_BGR2BGR565 || code == CV_RGB2BGR565 ||\r
+ // code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5, out, dcn,\r
+ // code == CV_BGR2BGR565 || code == CV_BGR2BGR555 ||\r
+ // code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2,\r
+ // stream);\r
// break;\r
\r
case CV_BGR2GRAY: case CV_BGRA2GRAY: case CV_RGB2GRAY: case CV_RGBA2GRAY:\r
nppSafeCall( nppiRGBToYCbCr_8u_C3R(src.ptr<Npp8u>(), src.step, out.ptr<Npp8u>(), out.step, nppsz) );\r
{\r
static int coeffs[] = {0, 2, 1};\r
- improc::swapChannels_gpu((DevMem2D)out, (DevMem2D)out, 3, coeffs, 0);\r
+ improc::swapChannels_gpu_8u(out, out, 3, coeffs, 0);\r
}\r
break;\r
\r
{\r
static int coeffs[] = {0, 2, 1};\r
GpuMat src1(src.size(), src.type());\r
- improc::swapChannels_gpu((DevMem2D)src, (DevMem2D)src1, 3, coeffs, 0);\r
+ improc::swapChannels_gpu_8u(src, src1, 3, coeffs, 0);\r
nppSafeCall( nppiYCbCrToRGB_8u_C3R(src1.ptr<Npp8u>(), src1.step, out.ptr<Npp8u>(), out.step, nppsz) ); \r
} \r
break;\r
-/*M///////////////////////////////////////////////////////////////////////////////////////
-//
-// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
-//
-// By downloading, copying, installing or using the software you agree to this license.
-// If you do not agree to this license, do not download, install,
-// copy or use the software.
-//
-//
-// Intel License Agreement
-// For Open Source Computer Vision Library
-//
-// Copyright (C) 2000, Intel Corporation, all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// Redistribution and use in source and binary forms, with or without modification,
-// are permitted provided that the following conditions are met:
-//
-// * Redistribution's of source code must retain the above copyright notice,
-// this list of conditions and the following disclaimer.
-//
-// * Redistribution's in binary form must reproduce the above copyright notice,
-// this list of conditions and the following disclaimer in the documentation
-// and/or other materials provided with the distribution.
-//
-// * The name of Intel Corporation may not be used to endorse or promote products
-// derived from this software without specific prior written permission.
-//
-// This software is provided by the copyright holders and contributors "as is" and
-// any express or implied warranties, including, but not limited to, the implied
-// warranties of merchantability and fitness for a particular purpose are disclaimed.
-// In no event shall the Intel Corporation or contributors be liable for any direct,
-// indirect, incidental, special, exemplary, or consequential damages
-// (including, but not limited to, procurement of substitute goods or services;
-// loss of use, data, or profits; or business interruption) however caused
-// and on any theory of liability, whether in contract, strict liability,
-// or tort (including negligence or otherwise) arising in any way out of
-// the use of this software, even if advised of the possibility of such damage.
-//
-//M*/
-
-#include <iostream>
-#include <cmath>
-#include <limits>
-#include "gputest.hpp"
-#include "opencv2/core/core.hpp"
-#include "opencv2/imgproc/imgproc.hpp"
-#include "opencv2/highgui/highgui.hpp"
-
-using namespace cv;
-using namespace std;
-using namespace gpu;
-
-class CV_GpuArithmTest : public CvTest
-{
-public:
- CV_GpuArithmTest(const char* test_name, const char* test_funcs) : CvTest(test_name, test_funcs) {}
- virtual ~CV_GpuArithmTest() {}
-
-protected:
- void run(int);
-
- int test(int type);
-
- virtual int test(const Mat& mat1, const Mat& mat2) = 0;
-
- int CheckNorm(const Mat& m1, const Mat& m2);
- int CheckNorm(const Scalar& s1, const Scalar& s2);
- int CheckNorm(double d1, double d2);
-};
-
-int CV_GpuArithmTest::test(int type)
-{
- cv::Size sz(200, 200);
- cv::Mat mat1(sz, type), mat2(sz, type);
- cv::RNG rng(*ts->get_rng());
- rng.fill(mat1, cv::RNG::UNIFORM, cv::Scalar::all(10), cv::Scalar::all(100));
- rng.fill(mat2, cv::RNG::UNIFORM, cv::Scalar::all(10), cv::Scalar::all(100));
-
- return test(mat1, mat2);
-}
-
-int CV_GpuArithmTest::CheckNorm(const Mat& m1, const Mat& m2)
-{
- double ret = norm(m1, m2, NORM_INF);
-
- if (ret < std::numeric_limits<double>::epsilon())
- return CvTS::OK;
-
- ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
- return CvTS::FAIL_GENERIC;
-}
-
-int CV_GpuArithmTest::CheckNorm(const Scalar& s1, const Scalar& s2)
-{
- double ret0 = CheckNorm(s1[0], s2[0]), ret1 = CheckNorm(s1[1], s2[1]), ret2 = CheckNorm(s1[2], s2[2]), ret3 = CheckNorm(s1[3], s2[3]);
-
- return (ret0 == CvTS::OK && ret1 == CvTS::OK && ret2 == CvTS::OK && ret3 == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC;
-}
-
-int CV_GpuArithmTest::CheckNorm(double d1, double d2)
-{
- double ret = ::fabs(d1 - d2);
-
- if (ret < std::numeric_limits<double>::epsilon())
- return CvTS::OK;
-
- ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
- return CvTS::FAIL_GENERIC;
-}
-
-void CV_GpuArithmTest::run( int )
-{
- int testResult = CvTS::OK;
- try
- {
- const int types[] = {CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1};
- const char* type_names[] = {"CV_8UC1", "CV_8UC3", "CV_8UC4", "CV_32FC1"};
- const int type_count = sizeof(types)/sizeof(types[0]);
-
- //run tests
- for (int t = 0; t < type_count; ++t)
- {
- ts->printf(CvTS::LOG, "========Start test %s========\n", type_names[t]);
-
- if (CvTS::OK == test(types[t]))
- ts->printf(CvTS::LOG, "SUCCESS\n");
- else
- {
- ts->printf(CvTS::LOG, "FAIL\n");
- testResult = CvTS::FAIL_MISMATCH;
- }
- }
-
- ///!!! author, please remove commented code if loop above is equivalent.
-
-
- /*ts->printf(CvTS::LOG, "\n========Start test 8UC1========\n");
- if (test(CV_8UC1) == CvTS::OK)
- ts->printf(CvTS::LOG, "\nSUCCESS\n");
- else
- {
- ts->printf(CvTS::LOG, "\nFAIL\n");
- testResult = CvTS::FAIL_GENERIC;
- }
-
- ts->printf(CvTS::LOG, "\n========Start test 8UC3========\n");
- if (test(CV_8UC3) == CvTS::OK)
- ts->printf(CvTS::LOG, "\nSUCCESS\n");
- else
- {
- ts->printf(CvTS::LOG, "\nFAIL\n");
- testResult = CvTS::FAIL_GENERIC;
- }
-
- ts->printf(CvTS::LOG, "\n========Start test 8UC4========\n");
- if (test(CV_8UC4) == CvTS::OK)
- ts->printf(CvTS::LOG, "\nSUCCESS\n");
- else
- {
- ts->printf(CvTS::LOG, "\nFAIL\n");
- testResult = CvTS::FAIL_GENERIC;
- }
-
- ts->printf(CvTS::LOG, "\n========Start test 32FC1========\n");
- if (test(CV_32FC1) == CvTS::OK)
- ts->printf(CvTS::LOG, "\nSUCCESS\n");
- else
- {
- ts->printf(CvTS::LOG, "\nFAIL\n");
- testResult = CvTS::FAIL_GENERIC;
- }*/
- }
- catch(const cv::Exception& e)
- {
- if (!check_and_treat_gpu_exception(e, ts))
- throw;
- return;
- }
-
- ts->set_failed_test_info(testResult);
-}
-
-////////////////////////////////////////////////////////////////////////////////
-// Add
-
-struct CV_GpuNppImageAddTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageAddTest() : CV_GpuArithmTest( "GPU-NppImageAdd", "add" ) {}
-
- virtual int test(const Mat& mat1, const Mat& mat2)
- {
- if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- cv::Mat cpuRes;
- cv::add(mat1, mat2, cpuRes);
-
- GpuMat gpu1(mat1);
- GpuMat gpu2(mat2);
- GpuMat gpuRes;
- cv::gpu::add(gpu1, gpu2, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// Sub
-struct CV_GpuNppImageSubtractTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageSubtractTest() : CV_GpuArithmTest( "GPU-NppImageSubtract", "subtract" ) {}
-
- int test( const Mat& mat1, const Mat& mat2 )
- {
- if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- cv::Mat cpuRes;
- cv::subtract(mat1, mat2, cpuRes);
-
- GpuMat gpu1(mat1);
- GpuMat gpu2(mat2);
- GpuMat gpuRes;
- cv::gpu::subtract(gpu1, gpu2, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// multiply
-struct CV_GpuNppImageMultiplyTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageMultiplyTest() : CV_GpuArithmTest( "GPU-NppImageMultiply", "multiply" ) {}
-
- int test( const Mat& mat1, const Mat& mat2 )
- {
- if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- cv::Mat cpuRes;
- cv::multiply(mat1, mat2, cpuRes);
-
- GpuMat gpu1(mat1);
- GpuMat gpu2(mat2);
- GpuMat gpuRes;
- cv::gpu::multiply(gpu1, gpu2, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// divide
-struct CV_GpuNppImageDivideTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageDivideTest() : CV_GpuArithmTest( "GPU-NppImageDivide", "divide" ) {}
-
- int test( const Mat& mat1, const Mat& mat2 )
- {
- if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- cv::Mat cpuRes;
- cv::divide(mat1, mat2, cpuRes);
-
- GpuMat gpu1(mat1);
- GpuMat gpu2(mat2);
- GpuMat gpuRes;
- cv::gpu::divide(gpu1, gpu2, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// transpose
-struct CV_GpuNppImageTransposeTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageTransposeTest() : CV_GpuArithmTest( "GPU-NppImageTranspose", "transpose" ) {}
-
- int test( const Mat& mat1, const Mat& )
- {
- if (mat1.type() != CV_8UC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- cv::Mat cpuRes;
- cv::transpose(mat1, cpuRes);
-
- GpuMat gpu1(mat1);
- GpuMat gpuRes;
- cv::gpu::transpose(gpu1, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// absdiff
-struct CV_GpuNppImageAbsdiffTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageAbsdiffTest() : CV_GpuArithmTest( "GPU-NppImageAbsdiff", "absdiff" ) {}
-
- int test( const Mat& mat1, const Mat& mat2 )
- {
- if (mat1.type() != CV_8UC1 && mat1.type() != CV_32FC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- cv::Mat cpuRes;
- cv::absdiff(mat1, mat2, cpuRes);
-
- GpuMat gpu1(mat1);
- GpuMat gpu2(mat2);
- GpuMat gpuRes;
- cv::gpu::absdiff(gpu1, gpu2, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// compare
-struct CV_GpuNppImageCompareTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageCompareTest() : CV_GpuArithmTest( "GPU-NppImageCompare", "compare" ) {}
-
- int test( const Mat& mat1, const Mat& mat2 )
- {
- if (mat1.type() != CV_32FC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- int cmp_codes[] = {CMP_EQ, CMP_GT, CMP_GE, CMP_LT, CMP_LE, CMP_NE};
- const char* cmp_str[] = {"CMP_EQ", "CMP_GT", "CMP_GE", "CMP_LT", "CMP_LE", "CMP_NE"};
- int cmp_num = sizeof(cmp_codes) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < cmp_num; ++i)
- {
- ts->printf(CvTS::LOG, "\nCompare operation: %s\n", cmp_str[i]);
-
- cv::Mat cpuRes;
- cv::compare(mat1, mat2, cpuRes, cmp_codes[i]);
-
- GpuMat gpu1(mat1);
- GpuMat gpu2(mat2);
- GpuMat gpuRes;
- cv::gpu::compare(gpu1, gpu2, gpuRes, cmp_codes[i]);
-
- if (CheckNorm(cpuRes, gpuRes) != CvTS::OK)
- test_res = CvTS::FAIL_GENERIC;
- }
-
- return test_res;
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// meanStdDev
-struct CV_GpuNppImageMeanStdDevTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageMeanStdDevTest() : CV_GpuArithmTest( "GPU-NppImageMeanStdDev", "meanStdDev" ) {}
-
- int test( const Mat& mat1, const Mat& )
- {
- if (mat1.type() != CV_8UC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- Scalar cpumean;
- Scalar cpustddev;
- cv::meanStdDev(mat1, cpumean, cpustddev);
-
- GpuMat gpu1(mat1);
- Scalar gpumean;
- Scalar gpustddev;
- cv::gpu::meanStdDev(gpu1, gpumean, gpustddev);
-
- int test_res = CvTS::OK;
-
- if (CheckNorm(cpumean, gpumean) != CvTS::OK)
- {
- ts->printf(CvTS::LOG, "\nMean FAILED\n");
- test_res = CvTS::FAIL_GENERIC;
- }
-
- if (CheckNorm(cpustddev, gpustddev) != CvTS::OK)
- {
- ts->printf(CvTS::LOG, "\nStdDev FAILED\n");
- test_res = CvTS::FAIL_GENERIC;
- }
-
- return test_res;
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// norm
-struct CV_GpuNppImageNormTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageNormTest() : CV_GpuArithmTest( "GPU-NppImageNorm", "norm" ) {}
-
- int test( const Mat& mat1, const Mat& mat2 )
- {
- if (mat1.type() != CV_8UC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- int norms[] = {NORM_INF, NORM_L1, NORM_L2};
- const char* norms_str[] = {"NORM_INF", "NORM_L1", "NORM_L2"};
- int norms_num = sizeof(norms) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < norms_num; ++i)
- {
- ts->printf(CvTS::LOG, "\nNorm type: %s\n", norms_str[i]);
-
- double cpu_norm = cv::norm(mat1, mat2, norms[i]);
-
- GpuMat gpu1(mat1);
- GpuMat gpu2(mat2);
- double gpu_norm = cv::gpu::norm(gpu1, gpu2, norms[i]);
-
- if (CheckNorm(cpu_norm, gpu_norm) != CvTS::OK)
- test_res = CvTS::FAIL_GENERIC;
- }
-
- return test_res;
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// flip
-struct CV_GpuNppImageFlipTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageFlipTest() : CV_GpuArithmTest( "GPU-NppImageFlip", "flip" ) {}
-
- int test( const Mat& mat1, const Mat& )
- {
- if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- int flip_codes[] = {0, 1, -1};
- const char* flip_axis[] = {"X", "Y", "Both"};
- int flip_codes_num = sizeof(flip_codes) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < flip_codes_num; ++i)
- {
- ts->printf(CvTS::LOG, "\nFlip Axis: %s\n", flip_axis[i]);
-
- Mat cpu_res;
- cv::flip(mat1, cpu_res, flip_codes[i]);
-
- GpuMat gpu1(mat1);
- GpuMat gpu_res;
- cv::gpu::flip(gpu1, gpu_res, flip_codes[i]);
-
- if (CheckNorm(cpu_res, gpu_res) != CvTS::OK)
- test_res = CvTS::FAIL_GENERIC;
- }
-
- return test_res;
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// sum
-struct CV_GpuNppImageSumTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageSumTest() : CV_GpuArithmTest( "GPU-NppImageSum", "sum" ) {}
-
- int test( const Mat& mat1, const Mat& )
- {
- if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- Scalar cpures = cv::sum(mat1);
-
- GpuMat gpu1(mat1);
- Scalar gpures = cv::gpu::sum(gpu1);
-
- return CheckNorm(cpures, gpures);
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// minNax
-struct CV_GpuNppImageMinNaxTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageMinNaxTest() : CV_GpuArithmTest( "GPU-NppImageMinNax", "minNax" ) {}
-
- int test( const Mat& mat1, const Mat& )
- {
- if (mat1.type() != CV_8UC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- double cpumin, cpumax;
- cv::minMaxLoc(mat1, &cpumin, &cpumax);
-
- GpuMat gpu1(mat1);
- double gpumin, gpumax;
- cv::gpu::minMax(gpu1, &gpumin, &gpumax);
-
- return (CheckNorm(cpumin, gpumin) == CvTS::OK && CheckNorm(cpumax, gpumax) == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC;
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// LUT
-struct CV_GpuNppImageLUTTest : public CV_GpuArithmTest
-{
- CV_GpuNppImageLUTTest() : CV_GpuArithmTest( "GPU-NppImageLUT", "LUT" ) {}
-
- int test( const Mat& mat1, const Mat& )
- {
- if (mat1.type() != CV_8UC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- cv::Mat lut(1, 256, CV_32SC1);
- cv::RNG rng(*ts->get_rng());
- rng.fill(lut, cv::RNG::UNIFORM, cv::Scalar::all(100), cv::Scalar::all(200));
-
- cv::Mat cpuRes;
- cv::LUT(mat1, lut, cpuRes);
- cpuRes.convertTo(cpuRes, CV_8U);
-
- cv::gpu::GpuMat gpuRes;
- cv::gpu::LUT(GpuMat(mat1), lut, gpuRes);
-
- return CheckNorm(cpuRes, gpuRes);
- }
-};
-
-/////////////////////////////////////////////////////////////////////////////
-/////////////////// tests registration /////////////////////////////////////
-/////////////////////////////////////////////////////////////////////////////
-
-// If we comment some tests, we may foget/miss to uncomment it after.
-// Placing all test definitions in one place
-// makes us know about what tests are commented.
-
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// Intel License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000, Intel Corporation, all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of Intel Corporation may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include <iostream>\r
+#include <cmath>\r
+#include <limits>\r
+#include "gputest.hpp"\r
+#include "opencv2/core/core.hpp"\r
+#include "opencv2/imgproc/imgproc.hpp"\r
+#include "opencv2/highgui/highgui.hpp"\r
+\r
+using namespace cv;\r
+using namespace std;\r
+using namespace gpu;\r
+\r
+class CV_GpuArithmTest : public CvTest\r
+{\r
+public:\r
+ CV_GpuArithmTest(const char* test_name, const char* test_funcs) : CvTest(test_name, test_funcs) {}\r
+ virtual ~CV_GpuArithmTest() {}\r
+\r
+protected:\r
+ void run(int);\r
+ \r
+ int test(int type);\r
+\r
+ virtual int test(const Mat& mat1, const Mat& mat2) = 0;\r
+\r
+ int CheckNorm(const Mat& m1, const Mat& m2);\r
+ int CheckNorm(const Scalar& s1, const Scalar& s2);\r
+ int CheckNorm(double d1, double d2);\r
+};\r
+\r
+int CV_GpuArithmTest::test(int type)\r
+{\r
+ cv::Size sz(200, 200);\r
+ cv::Mat mat1(sz, type), mat2(sz, type);\r
+ cv::RNG rng(*ts->get_rng());\r
+ rng.fill(mat1, cv::RNG::UNIFORM, cv::Scalar::all(10), cv::Scalar::all(100));\r
+ rng.fill(mat2, cv::RNG::UNIFORM, cv::Scalar::all(10), cv::Scalar::all(100));\r
+\r
+ return test(mat1, mat2);\r
+}\r
+\r
+int CV_GpuArithmTest::CheckNorm(const Mat& m1, const Mat& m2)\r
+{\r
+ double ret = norm(m1, m2, NORM_INF);\r
+\r
+ if (ret < std::numeric_limits<double>::epsilon())\r
+ return CvTS::OK;\r
+\r
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);\r
+ return CvTS::FAIL_GENERIC;\r
+}\r
+\r
+int CV_GpuArithmTest::CheckNorm(const Scalar& s1, const Scalar& s2)\r
+{\r
+ double ret0 = CheckNorm(s1[0], s2[0]), ret1 = CheckNorm(s1[1], s2[1]), ret2 = CheckNorm(s1[2], s2[2]), ret3 = CheckNorm(s1[3], s2[3]);\r
+\r
+ return (ret0 == CvTS::OK && ret1 == CvTS::OK && ret2 == CvTS::OK && ret3 == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC;\r
+}\r
+\r
+int CV_GpuArithmTest::CheckNorm(double d1, double d2)\r
+{\r
+ double ret = ::fabs(d1 - d2);\r
+\r
+ if (ret < std::numeric_limits<double>::epsilon())\r
+ return CvTS::OK;\r
+\r
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);\r
+ return CvTS::FAIL_GENERIC;\r
+}\r
+\r
+void CV_GpuArithmTest::run( int )\r
+{\r
+ int testResult = CvTS::OK;\r
+ try\r
+ {\r
+ const int types[] = {CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1};\r
+ const char* type_names[] = {"CV_8UC1", "CV_8UC3", "CV_8UC4", "CV_32FC1"};\r
+ const int type_count = sizeof(types)/sizeof(types[0]);\r
+\r
+ //run tests\r
+ for (int t = 0; t < type_count; ++t)\r
+ {\r
+ ts->printf(CvTS::LOG, "========Start test %s========\n", type_names[t]);\r
+\r
+ if (CvTS::OK == test(types[t]))\r
+ ts->printf(CvTS::LOG, "SUCCESS\n");\r
+ else\r
+ {\r
+ ts->printf(CvTS::LOG, "FAIL\n");\r
+ testResult = CvTS::FAIL_MISMATCH;\r
+ } \r
+ }\r
+ }\r
+ catch(const cv::Exception& e)\r
+ {\r
+ if (!check_and_treat_gpu_exception(e, ts))\r
+ throw;\r
+ return;\r
+ }\r
+\r
+ ts->set_failed_test_info(testResult);\r
+}\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// Add\r
+\r
+struct CV_GpuNppImageAddTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageAddTest() : CV_GpuArithmTest( "GPU-NppImageAdd", "add" ) {}\r
+\r
+ virtual int test(const Mat& mat1, const Mat& mat2)\r
+ {\r
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::Mat cpuRes;\r
+ cv::add(mat1, mat2, cpuRes);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu2(mat2);\r
+ GpuMat gpuRes;\r
+ cv::gpu::add(gpu1, gpu2, gpuRes);\r
+\r
+ return CheckNorm(cpuRes, gpuRes);\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// Sub\r
+struct CV_GpuNppImageSubtractTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageSubtractTest() : CV_GpuArithmTest( "GPU-NppImageSubtract", "subtract" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& mat2 )\r
+ {\r
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::Mat cpuRes;\r
+ cv::subtract(mat1, mat2, cpuRes);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu2(mat2);\r
+ GpuMat gpuRes;\r
+ cv::gpu::subtract(gpu1, gpu2, gpuRes); \r
+\r
+ return CheckNorm(cpuRes, gpuRes);\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// multiply\r
+struct CV_GpuNppImageMultiplyTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageMultiplyTest() : CV_GpuArithmTest( "GPU-NppImageMultiply", "multiply" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& mat2 )\r
+ {\r
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::Mat cpuRes;\r
+ cv::multiply(mat1, mat2, cpuRes);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu2(mat2);\r
+ GpuMat gpuRes;\r
+ cv::gpu::multiply(gpu1, gpu2, gpuRes);\r
+\r
+ return CheckNorm(cpuRes, gpuRes);\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// divide\r
+struct CV_GpuNppImageDivideTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageDivideTest() : CV_GpuArithmTest( "GPU-NppImageDivide", "divide" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& mat2 )\r
+ {\r
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4 && mat1.type() != CV_32FC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::Mat cpuRes;\r
+ cv::divide(mat1, mat2, cpuRes);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu2(mat2);\r
+ GpuMat gpuRes;\r
+ cv::gpu::divide(gpu1, gpu2, gpuRes);\r
+\r
+ return CheckNorm(cpuRes, gpuRes);\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// transpose\r
+struct CV_GpuNppImageTransposeTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageTransposeTest() : CV_GpuArithmTest( "GPU-NppImageTranspose", "transpose" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& )\r
+ {\r
+ if (mat1.type() != CV_8UC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::Mat cpuRes;\r
+ cv::transpose(mat1, cpuRes);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpuRes;\r
+ cv::gpu::transpose(gpu1, gpuRes);\r
+\r
+ return CheckNorm(cpuRes, gpuRes);\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// absdiff\r
+struct CV_GpuNppImageAbsdiffTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageAbsdiffTest() : CV_GpuArithmTest( "GPU-NppImageAbsdiff", "absdiff" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& mat2 )\r
+ {\r
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_32FC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::Mat cpuRes;\r
+ cv::absdiff(mat1, mat2, cpuRes);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu2(mat2);\r
+ GpuMat gpuRes;\r
+ cv::gpu::absdiff(gpu1, gpu2, gpuRes);\r
+\r
+ return CheckNorm(cpuRes, gpuRes);\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// compare\r
+struct CV_GpuNppImageCompareTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageCompareTest() : CV_GpuArithmTest( "GPU-NppImageCompare", "compare" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& mat2 )\r
+ {\r
+ if (mat1.type() != CV_32FC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ int cmp_codes[] = {CMP_EQ, CMP_GT, CMP_GE, CMP_LT, CMP_LE, CMP_NE};\r
+ const char* cmp_str[] = {"CMP_EQ", "CMP_GT", "CMP_GE", "CMP_LT", "CMP_LE", "CMP_NE"};\r
+ int cmp_num = sizeof(cmp_codes) / sizeof(int);\r
+\r
+ int test_res = CvTS::OK;\r
+\r
+ for (int i = 0; i < cmp_num; ++i)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nCompare operation: %s\n", cmp_str[i]);\r
+\r
+ cv::Mat cpuRes;\r
+ cv::compare(mat1, mat2, cpuRes, cmp_codes[i]);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu2(mat2);\r
+ GpuMat gpuRes;\r
+ cv::gpu::compare(gpu1, gpu2, gpuRes, cmp_codes[i]);\r
+\r
+ if (CheckNorm(cpuRes, gpuRes) != CvTS::OK)\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ return test_res;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// meanStdDev\r
+struct CV_GpuNppImageMeanStdDevTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageMeanStdDevTest() : CV_GpuArithmTest( "GPU-NppImageMeanStdDev", "meanStdDev" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& )\r
+ {\r
+ if (mat1.type() != CV_8UC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ Scalar cpumean; \r
+ Scalar cpustddev;\r
+ cv::meanStdDev(mat1, cpumean, cpustddev);\r
+\r
+ GpuMat gpu1(mat1);\r
+ Scalar gpumean; \r
+ Scalar gpustddev;\r
+ cv::gpu::meanStdDev(gpu1, gpumean, gpustddev);\r
+\r
+ int test_res = CvTS::OK;\r
+\r
+ if (CheckNorm(cpumean, gpumean) != CvTS::OK)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nMean FAILED\n");\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ if (CheckNorm(cpustddev, gpustddev) != CvTS::OK)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nStdDev FAILED\n");\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ return test_res;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// norm\r
+struct CV_GpuNppImageNormTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageNormTest() : CV_GpuArithmTest( "GPU-NppImageNorm", "norm" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& mat2 )\r
+ {\r
+ if (mat1.type() != CV_8UC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ int norms[] = {NORM_INF, NORM_L1, NORM_L2};\r
+ const char* norms_str[] = {"NORM_INF", "NORM_L1", "NORM_L2"};\r
+ int norms_num = sizeof(norms) / sizeof(int);\r
+\r
+ int test_res = CvTS::OK;\r
+\r
+ for (int i = 0; i < norms_num; ++i)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nNorm type: %s\n", norms_str[i]);\r
+\r
+ double cpu_norm = cv::norm(mat1, mat2, norms[i]);\r
+\r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu2(mat2);\r
+ double gpu_norm = cv::gpu::norm(gpu1, gpu2, norms[i]);\r
+\r
+ if (CheckNorm(cpu_norm, gpu_norm) != CvTS::OK)\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ return test_res;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// flip\r
+struct CV_GpuNppImageFlipTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageFlipTest() : CV_GpuArithmTest( "GPU-NppImageFlip", "flip" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& )\r
+ {\r
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ int flip_codes[] = {0, 1, -1};\r
+ const char* flip_axis[] = {"X", "Y", "Both"};\r
+ int flip_codes_num = sizeof(flip_codes) / sizeof(int);\r
+\r
+ int test_res = CvTS::OK;\r
+\r
+ for (int i = 0; i < flip_codes_num; ++i)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nFlip Axis: %s\n", flip_axis[i]);\r
+\r
+ Mat cpu_res;\r
+ cv::flip(mat1, cpu_res, flip_codes[i]);\r
+ \r
+ GpuMat gpu1(mat1);\r
+ GpuMat gpu_res;\r
+ cv::gpu::flip(gpu1, gpu_res, flip_codes[i]);\r
+\r
+ if (CheckNorm(cpu_res, gpu_res) != CvTS::OK)\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ return test_res;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// sum\r
+struct CV_GpuNppImageSumTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageSumTest() : CV_GpuArithmTest( "GPU-NppImageSum", "sum" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& )\r
+ {\r
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC4)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ Scalar cpures = cv::sum(mat1);\r
+\r
+ GpuMat gpu1(mat1);\r
+ Scalar gpures = cv::gpu::sum(gpu1);\r
+\r
+ return CheckNorm(cpures, gpures);\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// minNax\r
+struct CV_GpuNppImageMinNaxTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageMinNaxTest() : CV_GpuArithmTest( "GPU-NppImageMinNax", "minNax" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& )\r
+ {\r
+ if (mat1.type() != CV_8UC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ double cpumin, cpumax;\r
+ cv::minMaxLoc(mat1, &cpumin, &cpumax);\r
+\r
+ GpuMat gpu1(mat1);\r
+ double gpumin, gpumax;\r
+ cv::gpu::minMax(gpu1, &gpumin, &gpumax);\r
+\r
+ return (CheckNorm(cpumin, gpumin) == CvTS::OK && CheckNorm(cpumax, gpumax) == CvTS::OK) ? CvTS::OK : CvTS::FAIL_GENERIC;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// LUT\r
+struct CV_GpuNppImageLUTTest : public CV_GpuArithmTest\r
+{\r
+ CV_GpuNppImageLUTTest() : CV_GpuArithmTest( "GPU-NppImageLUT", "LUT" ) {}\r
+\r
+ int test( const Mat& mat1, const Mat& )\r
+ {\r
+ if (mat1.type() != CV_8UC1 && mat1.type() != CV_8UC3)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::Mat lut(1, 256, CV_8UC1);\r
+ cv::RNG rng(*ts->get_rng());\r
+ rng.fill(lut, cv::RNG::UNIFORM, cv::Scalar::all(100), cv::Scalar::all(200));\r
+\r
+ cv::Mat cpuRes;\r
+ cv::LUT(mat1, lut, cpuRes);\r
+\r
+ cv::gpu::GpuMat gpuRes;\r
+ cv::gpu::LUT(GpuMat(mat1), lut, gpuRes);\r
+\r
+ return CheckNorm(cpuRes, gpuRes);\r
+ }\r
+};\r
+\r
+/////////////////////////////////////////////////////////////////////////////\r
+/////////////////// tests registration /////////////////////////////////////\r
+/////////////////////////////////////////////////////////////////////////////\r
+\r
+// If we comment some tests, we may foget/miss to uncomment it after.\r
+// Placing all test definitions in one place \r
+// makes us know about what tests are commented.\r
+\r
CV_GpuNppImageAddTest CV_GpuNppImageAdd_test;\r
CV_GpuNppImageSubtractTest CV_GpuNppImageSubtract_test;\r
CV_GpuNppImageMultiplyTest CV_GpuNppImageMultiply_test;\r
const char* blacklist[] =
{
"GPU-NppImageSum",
+ "GPU-MatOperatorAsyncCall",
+ //"GPU-NppErode",
+ //"GPU-NppDilate",
+ //"GPU-NppMorphologyEx",
+ //"GPU-NppImageDivide",
+ //"GPU-NppImageMeanStdDev",
+ //"GPU-NppImageMinNax",
+ //"GPU-NppImageResize",
+ //"GPU-NppImageWarpAffine",
+ //"GPU-NppImageWarpPerspective",
+ //"GPU-NppImageIntegral",
+ //"GPU-NppImageBlur",
0
};
-/*M///////////////////////////////////////////////////////////////////////////////////////
-//
-// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
-//
-// By downloading, copying, installing or using the software you agree to this license.
-// If you do not agree to this license, do not download, install,
-// copy or use the software.
-//
-//
-// Intel License Agreement
-// For Open Source Computer Vision Library
-//
-// Copyright (C) 2000, Intel Corporation, all rights reserved.
-// Third party copyrights are property of their respective owners.
-//
-// Redistribution and use in source and binary forms, with or without modification,
-// are permitted provided that the following conditions are met:
-//
-// * Redistribution's of source code must retain the above copyright notice,
-// this list of conditions and the following disclaimer.
-//
-// * Redistribution's in binary form must reproduce the above copyright notice,
-// this list of conditions and the following disclaimer in the documentation
-// and/or other materials provided with the distribution.
-//
-// * The name of Intel Corporation may not be used to endorse or promote products
-// derived from this software without specific prior written permission.
-//
-// This software is provided by the copyright holders and contributors "as is" and
-// any express or implied warranties, including, but not limited to, the implied
-// warranties of merchantability and fitness for a particular purpose are disclaimed.
-// In no event shall the Intel Corporation or contributors be liable for any direct,
-// indirect, incidental, special, exemplary, or consequential damages
-// (including, but not limited to, procurement of substitute goods or services;
-// loss of use, data, or profits; or business interruption) however caused
-// and on any theory of liability, whether in contract, strict liability,
-// or tort (including negligence or otherwise) arising in any way out of
-// the use of this software, even if advised of the possibility of such damage.
-//
-//M*/
-
-#include <iostream>
-#include <cmath>
-#include <limits>
-#include "gputest.hpp"
-#include "opencv2/core/core.hpp"
-#include "opencv2/imgproc/imgproc.hpp"
-#include "opencv2/highgui/highgui.hpp"
-
-using namespace cv;
-using namespace std;
-using namespace gpu;
-
-class CV_GpuImageProcTest : public CvTest
-{
-public:
- CV_GpuImageProcTest(const char* test_name, const char* test_funcs) : CvTest(test_name, test_funcs) {}
- virtual ~CV_GpuImageProcTest() {}
-
-protected:
- void run(int);
-
- int test8UC1 (const Mat& img);
- int test8UC4 (const Mat& img);
- int test32SC1(const Mat& img);
- int test32FC1(const Mat& img);
-
- virtual int test(const Mat& img) = 0;
-
- int CheckNorm(const Mat& m1, const Mat& m2);
-};
-
-
-int CV_GpuImageProcTest::test8UC1(const Mat& img)
-{
- cv::Mat img_C1;
- cvtColor(img, img_C1, CV_BGR2GRAY);
-
- return test(img_C1);
-}
-
-int CV_GpuImageProcTest::test8UC4(const Mat& img)
-{
- cv::Mat img_C4;
- cvtColor(img, img_C4, CV_BGR2BGRA);
-
- return test(img_C4);
-}
-
-int CV_GpuImageProcTest::test32SC1(const Mat& img)
-{
- cv::Mat img_C1;
- cvtColor(img, img_C1, CV_BGR2GRAY);
- img_C1.convertTo(img_C1, CV_32S);
-
- return test(img_C1);
-}
-
-int CV_GpuImageProcTest::test32FC1(const Mat& img)
-{
- cv::Mat temp, img_C1;
- img.convertTo(temp, CV_32F);
- cvtColor(temp, img_C1, CV_BGR2GRAY);
-
- return test(img_C1);
-}
-
-int CV_GpuImageProcTest::CheckNorm(const Mat& m1, const Mat& m2)
-{
- double ret = norm(m1, m2, NORM_INF);
-
- if (ret < std::numeric_limits<double>::epsilon())
- {
- return CvTS::OK;
- }
- else
- {
- ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
- return CvTS::FAIL_GENERIC;
- }
-}
-
-void CV_GpuImageProcTest::run( int )
-{
- //load image
- cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");
-
- if (img.empty())
- {
- ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
- return;
- }
-
- int testResult = CvTS::OK;
- try
- {
- //run tests
- ts->printf(CvTS::LOG, "\n========Start test 8UC1========\n");
- if (test8UC1(img) == CvTS::OK)
- ts->printf(CvTS::LOG, "\nSUCCESS\n");
- else
- {
- ts->printf(CvTS::LOG, "\nFAIL\n");
- testResult = CvTS::FAIL_GENERIC;
- }
-
- ts->printf(CvTS::LOG, "\n========Start test 8UC4========\n");
- if (test8UC4(img) == CvTS::OK)
- ts->printf(CvTS::LOG, "\nSUCCESS\n");
- else
- {
- ts->printf(CvTS::LOG, "\nFAIL\n");
- testResult = CvTS::FAIL_GENERIC;
- }
-
- ts->printf(CvTS::LOG, "\n========Start test 32SC1========\n");
- if (test32SC1(img) == CvTS::OK)
- ts->printf(CvTS::LOG, "\nSUCCESS\n");
- else
- {
- ts->printf(CvTS::LOG, "\nFAIL\n");
- testResult = CvTS::FAIL_GENERIC;
- }
-
- ts->printf(CvTS::LOG, "\n========Start test 32FC1========\n");
- if (test32FC1(img) == CvTS::OK)
- ts->printf(CvTS::LOG, "\nSUCCESS\n");
- else
- {
- ts->printf(CvTS::LOG, "\nFAIL\n");
- testResult = CvTS::FAIL_GENERIC;
- }
- }
- catch(const cv::Exception& e)
- {
- if (!check_and_treat_gpu_exception(e, ts))
- throw;
- return;
- }
-
- ts->set_failed_test_info(testResult);
-}
-
-////////////////////////////////////////////////////////////////////////////////
-// threshold
-struct CV_GpuNppImageThresholdTest : public CV_GpuImageProcTest
-{
-public:
- CV_GpuNppImageThresholdTest() : CV_GpuImageProcTest( "GPU-NppImageThreshold", "threshold" ) {}
-
- int test(const Mat& img)
- {
- if (img.type() != CV_32FC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- cv::RNG rng(*ts->get_rng());
- const double thresh = rng;
-
- cv::Mat cpuRes;
- cv::threshold(img, cpuRes, thresh, 0.0, THRESH_TRUNC);
-
- GpuMat gpu1(img);
- GpuMat gpuRes;
- cv::gpu::threshold(gpu1, gpuRes, thresh);
-
- return CheckNorm(cpuRes, gpuRes);
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// resize
-struct CV_GpuNppImageResizeTest : public CV_GpuImageProcTest
-{
- CV_GpuNppImageResizeTest() : CV_GpuImageProcTest( "GPU-NppImageResize", "resize" ) {}
- int test(const Mat& img)
- {
- if (img.type() != CV_8UC1 && img.type() != CV_8UC4)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- int interpolations[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_LANCZOS4};
- const char* interpolations_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_LANCZOS4"};
- int interpolations_num = sizeof(interpolations) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < interpolations_num; ++i)
- {
- ts->printf(CvTS::LOG, "\nInterpolation type: %s\n", interpolations_str[i]);
-
- Mat cpu_res;
- cv::resize(img, cpu_res, Size(), 0.5, 0.5, interpolations[i]);
-
- GpuMat gpu1(img), gpu_res;
- cv::gpu::resize(gpu1, gpu_res, Size(), 0.5, 0.5, interpolations[i]);
-
- if (CheckNorm(cpu_res, gpu_res) != CvTS::OK)
- test_res = CvTS::FAIL_GENERIC;
- }
-
- return test_res;
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// copyMakeBorder
-struct CV_GpuNppImageCopyMakeBorderTest : public CV_GpuImageProcTest
-{
- CV_GpuNppImageCopyMakeBorderTest() : CV_GpuImageProcTest( "GPU-NppImageCopyMakeBorder", "copyMakeBorder" ) {}
-
- int test(const Mat& img)
- {
- if (img.type() != CV_8UC1 && img.type() != CV_8UC4 && img.type() != CV_32SC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- cv::RNG rng(*ts->get_rng());
- int top = rng.uniform(1, 10);
- int botton = rng.uniform(1, 10);
- int left = rng.uniform(1, 10);
- int right = rng.uniform(1, 10);
- cv::Scalar val(rng.uniform(0, 255), rng.uniform(0, 255), rng.uniform(0, 255), rng.uniform(0, 255));
-
- Mat cpudst;
- cv::copyMakeBorder(img, cpudst, top, botton, left, right, BORDER_CONSTANT, val);
-
- GpuMat gpu1(img);
- GpuMat gpudst;
- cv::gpu::copyMakeBorder(gpu1, gpudst, top, botton, left, right, val);
-
- return CheckNorm(cpudst, gpudst);
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// warpAffine
-struct CV_GpuNppImageWarpAffineTest : public CV_GpuImageProcTest
-{
- CV_GpuNppImageWarpAffineTest() : CV_GpuImageProcTest( "GPU-NppImageWarpAffine", "warpAffine" ) {}
-
- int test(const Mat& img)
- {
- if (img.type() == CV_32SC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- static const double coeffs[2][3] =
- {
- {cos(3.14 / 6), -sin(3.14 / 6), 100.0},
- {sin(3.14 / 6), cos(3.14 / 6), -100.0}
- };
- Mat M(2, 3, CV_64F, (void*)coeffs);
-
- int flags[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_NEAREST | WARP_INVERSE_MAP, INTER_LINEAR | WARP_INVERSE_MAP, INTER_CUBIC | WARP_INVERSE_MAP};
- const char* flags_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_NEAREST | WARP_INVERSE_MAP", "INTER_LINEAR | WARP_INVERSE_MAP", "INTER_CUBIC | WARP_INVERSE_MAP"};
- int flags_num = sizeof(flags) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < flags_num; ++i)
- {
- ts->printf(CvTS::LOG, "\nFlags: %s\n", flags_str[i]);
-
- Mat cpudst;
- cv::warpAffine(img, cpudst, M, img.size(), flags[i]);
-
- GpuMat gpu1(img);
- GpuMat gpudst;
- cv::gpu::warpAffine(gpu1, gpudst, M, gpu1.size(), flags[i]);
-
- if (CheckNorm(cpudst, gpudst) != CvTS::OK)
- test_res = CvTS::FAIL_GENERIC;
- }
-
- return test_res;
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// warpPerspective
-struct CV_GpuNppImageWarpPerspectiveTest : public CV_GpuImageProcTest
-{
- CV_GpuNppImageWarpPerspectiveTest() : CV_GpuImageProcTest( "GPU-NppImageWarpPerspective", "warpPerspective" ) {}
-
-
- int test(const Mat& img)
- {
- if (img.type() == CV_32SC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- static const double coeffs[3][3] =
- {
- {cos(3.14 / 6), -sin(3.14 / 6), 100.0},
- {sin(3.14 / 6), cos(3.14 / 6), -100.0},
- {0.0, 0.0, 1.0}
- };
- Mat M(3, 3, CV_64F, (void*)coeffs);
-
- int flags[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_NEAREST | WARP_INVERSE_MAP, INTER_LINEAR | WARP_INVERSE_MAP, INTER_CUBIC | WARP_INVERSE_MAP};
- const char* flags_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_NEAREST | WARP_INVERSE_MAP", "INTER_LINEAR | WARP_INVERSE_MAP", "INTER_CUBIC | WARP_INVERSE_MAP"};
- int flags_num = sizeof(flags) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < flags_num; ++i)
- {
- ts->printf(CvTS::LOG, "\nFlags: %s\n", flags_str[i]);
-
- Mat cpudst;
- cv::warpPerspective(img, cpudst, M, img.size(), flags[i]);
-
- GpuMat gpu1(img);
- GpuMat gpudst;
- cv::gpu::warpPerspective(gpu1, gpudst, M, gpu1.size(), flags[i]);
-
- if (CheckNorm(cpudst, gpudst) != CvTS::OK)
- test_res = CvTS::FAIL_GENERIC;
- }
-
- return test_res;
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// integral
-struct CV_GpuNppImageIntegralTest : public CV_GpuImageProcTest
-{
- CV_GpuNppImageIntegralTest() : CV_GpuImageProcTest( "GPU-NppImageIntegral", "integral" ) {}
-
- int CV_GpuNppImageIntegralTest::test(const Mat& img)
- {
- if (img.type() != CV_8UC1)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- Mat cpusum, cpusqsum;
- cv::integral(img, cpusum, cpusqsum, CV_32S);
-
- GpuMat gpu1(img);
- GpuMat gpusum, gpusqsum;
- cv::gpu::integral(gpu1, gpusum, gpusqsum);
-
- gpusqsum.convertTo(gpusqsum, CV_64F);
-
- int test_res = CvTS::OK;
-
- if (CheckNorm(cpusum, gpusum) != CvTS::OK)
- {
- ts->printf(CvTS::LOG, "\nSum failed\n");
- test_res = CvTS::FAIL_GENERIC;
- }
- if (CheckNorm(cpusqsum, gpusqsum) != CvTS::OK)
- {
- ts->printf(CvTS::LOG, "\nSquared sum failed\n");
- test_res = CvTS::FAIL_GENERIC;
- }
-
- return test_res;
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// blur
-struct CV_GpuNppImageBlurTest : public CV_GpuImageProcTest
-{
- CV_GpuNppImageBlurTest() : CV_GpuImageProcTest( "GPU-NppImageBlur", "blur" ) {}
-
- int test(const Mat& img)
- {
- if (img.type() != CV_8UC1 && img.type() != CV_8UC4)
- {
- ts->printf(CvTS::LOG, "\nUnsupported type\n");
- return CvTS::OK;
- }
-
- int ksizes[] = {3, 5, 7};
- int ksizes_num = sizeof(ksizes) / sizeof(int);
-
- int test_res = CvTS::OK;
-
- for (int i = 0; i < ksizes_num; ++i)
- {
- ts->printf(CvTS::LOG, "\nksize = %d\n", ksizes[i]);
-
- Mat cpudst;
- cv::blur(img, cpudst, Size(ksizes[i], ksizes[i]));
-
- GpuMat gpu1(img);
- GpuMat gpudst;
- cv::gpu::blur(gpu1, gpudst, Size(ksizes[i], ksizes[i]));
-
- cv::Mat c;
- cv::absdiff(cpudst, gpudst, c);
-
- if (CheckNorm(cpudst, gpudst) != CvTS::OK)
- test_res = CvTS::FAIL_GENERIC;
- }
-
- return test_res;
- }
-};
-
-////////////////////////////////////////////////////////////////////////////////
-// cvtColor
-class CV_GpuCvtColorTest : public CvTest
-{
-public:
- CV_GpuCvtColorTest() : CvTest("GPU-NppCvtColor", "cvtColor") {}
- ~CV_GpuCvtColorTest() {};
-
-protected:
- void run(int);
-
- int CheckNorm(const Mat& m1, const Mat& m2);
-};
-
-
-int CV_GpuCvtColorTest::CheckNorm(const Mat& m1, const Mat& m2)
-{
- double ret = norm(m1, m2, NORM_INF);
-
- if (ret < std::numeric_limits<double>::epsilon())
- {
- return CvTS::OK;
- }
- else
- {
- ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);
- return CvTS::FAIL_GENERIC;
- }
-}
-
-void CV_GpuCvtColorTest::run( int )
-{
- //load image
- cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");
-
- if (img.empty())
- {
- ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
- return;
- }
-
- int testResult = CvTS::OK;
- cv::Mat cpuRes;
- cv::gpu::GpuMat gpuImg(img), gpuRes;
- try
- {
- //run tests
- int codes[] = {CV_BGR2RGB, CV_RGB2YCrCb, CV_YCrCb2RGB, CV_RGB2RGBA, CV_RGBA2BGRA, CV_BGRA2GRAY, CV_GRAY2RGB};
- const char* codes_str[] = {"CV_BGR2RGB", "CV_RGB2YCrCb", "CV_YCrCb2RGB", "CV_RGB2RGBA", "CV_RGBA2BGRA", "CV_BGRA2GRAY", "CV_GRAY2RGB"};
- int codes_num = sizeof(codes) / sizeof(int);
-
- for (int i = 0; i < codes_num; ++i)
- {
- ts->printf(CvTS::LOG, "\n%s\n", codes_str[i]);
-
- cv::cvtColor(img, cpuRes, codes[i]);
- cv::gpu::cvtColor(gpuImg, gpuRes, codes[i]);
-
- if (CheckNorm(cpuRes, gpuRes) == CvTS::OK)
- ts->printf(CvTS::LOG, "\nSUCCESS\n");
- else
- {
- ts->printf(CvTS::LOG, "\nFAIL\n");
- testResult = CvTS::FAIL_GENERIC;
- }
-
- img = cpuRes;
- gpuImg = gpuRes;
- }
- }
- catch(const cv::Exception& e)
- {
- if (!check_and_treat_gpu_exception(e, ts))
- throw;
- return;
- }
-
- ts->set_failed_test_info(testResult);
-}
-
-/////////////////////////////////////////////////////////////////////////////
-/////////////////// tests registration /////////////////////////////////////
-/////////////////////////////////////////////////////////////////////////////
-
-// If we comment some tests, we may foget/miss to uncomment it after.
-// Placing all test definitions in one place
-// makes us know about what tests are commented.
-
+/*M///////////////////////////////////////////////////////////////////////////////////////\r
+//\r
+// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.\r
+//\r
+// By downloading, copying, installing or using the software you agree to this license.\r
+// If you do not agree to this license, do not download, install,\r
+// copy or use the software.\r
+//\r
+//\r
+// Intel License Agreement\r
+// For Open Source Computer Vision Library\r
+//\r
+// Copyright (C) 2000, Intel Corporation, all rights reserved.\r
+// Third party copyrights are property of their respective owners.\r
+//\r
+// Redistribution and use in source and binary forms, with or without modification,\r
+// are permitted provided that the following conditions are met:\r
+//\r
+// * Redistribution's of source code must retain the above copyright notice,\r
+// this list of conditions and the following disclaimer.\r
+//\r
+// * Redistribution's in binary form must reproduce the above copyright notice,\r
+// this list of conditions and the following disclaimer in the documentation\r
+// and/or other materials provided with the distribution.\r
+//\r
+// * The name of Intel Corporation may not be used to endorse or promote products\r
+// derived from this software without specific prior written permission.\r
+//\r
+// This software is provided by the copyright holders and contributors "as is" and\r
+// any express or implied warranties, including, but not limited to, the implied\r
+// warranties of merchantability and fitness for a particular purpose are disclaimed.\r
+// In no event shall the Intel Corporation or contributors be liable for any direct,\r
+// indirect, incidental, special, exemplary, or consequential damages\r
+// (including, but not limited to, procurement of substitute goods or services;\r
+// loss of use, data, or profits; or business interruption) however caused\r
+// and on any theory of liability, whether in contract, strict liability,\r
+// or tort (including negligence or otherwise) arising in any way out of\r
+// the use of this software, even if advised of the possibility of such damage.\r
+//\r
+//M*/\r
+\r
+#include <iostream>\r
+#include <cmath>\r
+#include <limits>\r
+#include "gputest.hpp"\r
+#include "opencv2/core/core.hpp"\r
+#include "opencv2/imgproc/imgproc.hpp"\r
+#include "opencv2/highgui/highgui.hpp"\r
+\r
+using namespace cv;\r
+using namespace std;\r
+using namespace gpu;\r
+\r
+class CV_GpuImageProcTest : public CvTest\r
+{\r
+public:\r
+ CV_GpuImageProcTest(const char* test_name, const char* test_funcs) : CvTest(test_name, test_funcs) {}\r
+ virtual ~CV_GpuImageProcTest() {}\r
+\r
+protected:\r
+ void run(int);\r
+ \r
+ int test8UC1 (const Mat& img);\r
+ int test8UC4 (const Mat& img);\r
+ int test32SC1(const Mat& img);\r
+ int test32FC1(const Mat& img);\r
+\r
+ virtual int test(const Mat& img) = 0;\r
+\r
+ int CheckNorm(const Mat& m1, const Mat& m2);\r
+};\r
+\r
+\r
+int CV_GpuImageProcTest::test8UC1(const Mat& img)\r
+{\r
+ cv::Mat img_C1;\r
+ cvtColor(img, img_C1, CV_BGR2GRAY);\r
+\r
+ return test(img_C1);\r
+}\r
+\r
+int CV_GpuImageProcTest::test8UC4(const Mat& img)\r
+{\r
+ cv::Mat img_C4;\r
+ cvtColor(img, img_C4, CV_BGR2BGRA);\r
+\r
+ return test(img_C4);\r
+}\r
+\r
+int CV_GpuImageProcTest::test32SC1(const Mat& img)\r
+{\r
+ cv::Mat img_C1;\r
+ cvtColor(img, img_C1, CV_BGR2GRAY); \r
+ img_C1.convertTo(img_C1, CV_32S);\r
+\r
+ return test(img_C1);\r
+}\r
+\r
+int CV_GpuImageProcTest::test32FC1(const Mat& img)\r
+{\r
+ cv::Mat temp, img_C1;\r
+ img.convertTo(temp, CV_32F);\r
+ cvtColor(temp, img_C1, CV_BGR2GRAY);\r
+\r
+ return test(img_C1);\r
+}\r
+\r
+int CV_GpuImageProcTest::CheckNorm(const Mat& m1, const Mat& m2)\r
+{\r
+ double ret = norm(m1, m2, NORM_INF);\r
+\r
+ if (ret < std::numeric_limits<double>::epsilon())\r
+ {\r
+ return CvTS::OK;\r
+ }\r
+ else\r
+ {\r
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);\r
+ return CvTS::FAIL_GENERIC;\r
+ }\r
+}\r
+\r
+void CV_GpuImageProcTest::run( int )\r
+{\r
+ //load image\r
+ cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");\r
+\r
+ if (img.empty())\r
+ {\r
+ ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);\r
+ return;\r
+ }\r
+\r
+ int testResult = CvTS::OK;\r
+ try\r
+ {\r
+ //run tests\r
+ ts->printf(CvTS::LOG, "\n========Start test 8UC1========\n");\r
+ if (test8UC1(img) == CvTS::OK)\r
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");\r
+ else\r
+ {\r
+ ts->printf(CvTS::LOG, "\nFAIL\n");\r
+ testResult = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ ts->printf(CvTS::LOG, "\n========Start test 8UC4========\n");\r
+ if (test8UC4(img) == CvTS::OK)\r
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");\r
+ else\r
+ {\r
+ ts->printf(CvTS::LOG, "\nFAIL\n");\r
+ testResult = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ ts->printf(CvTS::LOG, "\n========Start test 32SC1========\n");\r
+ if (test32SC1(img) == CvTS::OK)\r
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");\r
+ else\r
+ {\r
+ ts->printf(CvTS::LOG, "\nFAIL\n");\r
+ testResult = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ ts->printf(CvTS::LOG, "\n========Start test 32FC1========\n");\r
+ if (test32FC1(img) == CvTS::OK)\r
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");\r
+ else\r
+ {\r
+ ts->printf(CvTS::LOG, "\nFAIL\n");\r
+ testResult = CvTS::FAIL_GENERIC;\r
+ }\r
+ }\r
+ catch(const cv::Exception& e)\r
+ {\r
+ if (!check_and_treat_gpu_exception(e, ts))\r
+ throw;\r
+ return;\r
+ }\r
+\r
+ ts->set_failed_test_info(testResult);\r
+}\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// threshold\r
+struct CV_GpuNppImageThresholdTest : public CV_GpuImageProcTest\r
+{\r
+public:\r
+ CV_GpuNppImageThresholdTest() : CV_GpuImageProcTest( "GPU-NppImageThreshold", "threshold" ) {}\r
+\r
+ int test(const Mat& img)\r
+ {\r
+ if (img.type() != CV_32FC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::RNG rng(*ts->get_rng());\r
+ const double thresh = rng;\r
+\r
+ cv::Mat cpuRes;\r
+ cv::threshold(img, cpuRes, thresh, 0.0, THRESH_TRUNC);\r
+\r
+ GpuMat gpu1(img);\r
+ GpuMat gpuRes;\r
+ cv::gpu::threshold(gpu1, gpuRes, thresh);\r
+\r
+ return CheckNorm(cpuRes, gpuRes);\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// resize\r
+struct CV_GpuNppImageResizeTest : public CV_GpuImageProcTest\r
+{\r
+ CV_GpuNppImageResizeTest() : CV_GpuImageProcTest( "GPU-NppImageResize", "resize" ) {}\r
+ int test(const Mat& img)\r
+ {\r
+ if (img.type() != CV_8UC1 && img.type() != CV_8UC4)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ int interpolations[] = {INTER_NEAREST, INTER_LINEAR, /*INTER_CUBIC,*/ /*INTER_LANCZOS4*/};\r
+ const char* interpolations_str[] = {"INTER_NEAREST", "INTER_LINEAR", /*"INTER_CUBIC",*/ /*"INTER_LANCZOS4"*/};\r
+ int interpolations_num = sizeof(interpolations) / sizeof(int);\r
+\r
+ int test_res = CvTS::OK;\r
+\r
+ for (int i = 0; i < interpolations_num; ++i)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nInterpolation type: %s\n", interpolations_str[i]);\r
+\r
+ Mat cpu_res;\r
+ cv::resize(img, cpu_res, Size(), 0.5, 0.5, interpolations[i]);\r
+\r
+ GpuMat gpu1(img), gpu_res;\r
+ cv::gpu::resize(gpu1, gpu_res, Size(), 0.5, 0.5, interpolations[i]);\r
+\r
+ if (CheckNorm(cpu_res, gpu_res) != CvTS::OK)\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ return test_res;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// copyMakeBorder\r
+struct CV_GpuNppImageCopyMakeBorderTest : public CV_GpuImageProcTest\r
+{\r
+ CV_GpuNppImageCopyMakeBorderTest() : CV_GpuImageProcTest( "GPU-NppImageCopyMakeBorder", "copyMakeBorder" ) {}\r
+\r
+ int test(const Mat& img)\r
+ {\r
+ if (img.type() != CV_8UC1 && img.type() != CV_8UC4 && img.type() != CV_32SC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ cv::RNG rng(*ts->get_rng());\r
+ int top = rng.uniform(1, 10);\r
+ int botton = rng.uniform(1, 10);\r
+ int left = rng.uniform(1, 10);\r
+ int right = rng.uniform(1, 10);\r
+ cv::Scalar val(rng.uniform(0, 255), rng.uniform(0, 255), rng.uniform(0, 255), rng.uniform(0, 255));\r
+\r
+ Mat cpudst;\r
+ cv::copyMakeBorder(img, cpudst, top, botton, left, right, BORDER_CONSTANT, val);\r
+\r
+ GpuMat gpu1(img);\r
+ GpuMat gpudst; \r
+ cv::gpu::copyMakeBorder(gpu1, gpudst, top, botton, left, right, val);\r
+\r
+ return CheckNorm(cpudst, gpudst);\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// warpAffine\r
+struct CV_GpuNppImageWarpAffineTest : public CV_GpuImageProcTest\r
+{\r
+ CV_GpuNppImageWarpAffineTest() : CV_GpuImageProcTest( "GPU-NppImageWarpAffine", "warpAffine" ) {}\r
+\r
+ int test(const Mat& img)\r
+ {\r
+ if (img.type() == CV_32SC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+ \r
+ static const double coeffs[2][3] = \r
+ { \r
+ {cos(3.14 / 6), -sin(3.14 / 6), 100.0}, \r
+ {sin(3.14 / 6), cos(3.14 / 6), -100.0}\r
+ };\r
+ Mat M(2, 3, CV_64F, (void*)coeffs);\r
+\r
+ int flags[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_NEAREST | WARP_INVERSE_MAP, INTER_LINEAR | WARP_INVERSE_MAP, INTER_CUBIC | WARP_INVERSE_MAP};\r
+ const char* flags_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_NEAREST | WARP_INVERSE_MAP", "INTER_LINEAR | WARP_INVERSE_MAP", "INTER_CUBIC | WARP_INVERSE_MAP"};\r
+ int flags_num = sizeof(flags) / sizeof(int);\r
+\r
+ int test_res = CvTS::OK;\r
+\r
+ for (int i = 0; i < flags_num; ++i)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nFlags: %s\n", flags_str[i]);\r
+\r
+ Mat cpudst;\r
+ cv::warpAffine(img, cpudst, M, img.size(), flags[i]);\r
+\r
+ GpuMat gpu1(img);\r
+ GpuMat gpudst;\r
+ cv::gpu::warpAffine(gpu1, gpudst, M, gpu1.size(), flags[i]);\r
+ \r
+ if (CheckNorm(cpudst, gpudst) != CvTS::OK)\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ return test_res;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// warpPerspective\r
+struct CV_GpuNppImageWarpPerspectiveTest : public CV_GpuImageProcTest\r
+{\r
+ CV_GpuNppImageWarpPerspectiveTest() : CV_GpuImageProcTest( "GPU-NppImageWarpPerspective", "warpPerspective" ) {}\r
+\r
+\r
+ int test(const Mat& img)\r
+ {\r
+ if (img.type() == CV_32SC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+ \r
+ static const double coeffs[3][3] = \r
+ {\r
+ {cos(3.14 / 6), -sin(3.14 / 6), 100.0}, \r
+ {sin(3.14 / 6), cos(3.14 / 6), -100.0}, \r
+ {0.0, 0.0, 1.0}\r
+ };\r
+ Mat M(3, 3, CV_64F, (void*)coeffs);\r
+\r
+ int flags[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_NEAREST | WARP_INVERSE_MAP, INTER_LINEAR | WARP_INVERSE_MAP, INTER_CUBIC | WARP_INVERSE_MAP};\r
+ const char* flags_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_NEAREST | WARP_INVERSE_MAP", "INTER_LINEAR | WARP_INVERSE_MAP", "INTER_CUBIC | WARP_INVERSE_MAP"};\r
+ int flags_num = sizeof(flags) / sizeof(int);\r
+\r
+ int test_res = CvTS::OK;\r
+\r
+ for (int i = 0; i < flags_num; ++i)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nFlags: %s\n", flags_str[i]);\r
+\r
+ Mat cpudst;\r
+ cv::warpPerspective(img, cpudst, M, img.size(), flags[i]);\r
+\r
+ GpuMat gpu1(img);\r
+ GpuMat gpudst;\r
+ cv::gpu::warpPerspective(gpu1, gpudst, M, gpu1.size(), flags[i]);\r
+ \r
+ if (CheckNorm(cpudst, gpudst) != CvTS::OK)\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ return test_res;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// integral\r
+struct CV_GpuNppImageIntegralTest : public CV_GpuImageProcTest\r
+{\r
+ CV_GpuNppImageIntegralTest() : CV_GpuImageProcTest( "GPU-NppImageIntegral", "integral" ) {}\r
+\r
+ int CV_GpuNppImageIntegralTest::test(const Mat& img)\r
+ {\r
+ if (img.type() != CV_8UC1)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ Mat cpusum, cpusqsum;\r
+ cv::integral(img, cpusum, cpusqsum, CV_32S);\r
+\r
+ GpuMat gpu1(img);\r
+ GpuMat gpusum, gpusqsum;\r
+ cv::gpu::integral(gpu1, gpusum, gpusqsum);\r
+\r
+ gpusqsum.convertTo(gpusqsum, CV_64F);\r
+\r
+ int test_res = CvTS::OK;\r
+\r
+ if (CheckNorm(cpusum, gpusum) != CvTS::OK)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nSum failed\n");\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+ if (CheckNorm(cpusqsum, gpusqsum) != CvTS::OK)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nSquared sum failed\n");\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ return test_res;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// blur\r
+struct CV_GpuNppImageBlurTest : public CV_GpuImageProcTest\r
+{\r
+ CV_GpuNppImageBlurTest() : CV_GpuImageProcTest( "GPU-NppImageBlur", "blur" ) {}\r
+\r
+ int test(const Mat& img)\r
+ {\r
+ if (img.type() != CV_8UC1 && img.type() != CV_8UC4)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nUnsupported type\n");\r
+ return CvTS::OK;\r
+ }\r
+\r
+ int ksizes[] = {3, 5, 7};\r
+ int ksizes_num = sizeof(ksizes) / sizeof(int);\r
+\r
+ int test_res = CvTS::OK;\r
+\r
+ for (int i = 0; i < ksizes_num; ++i)\r
+ {\r
+ ts->printf(CvTS::LOG, "\nksize = %d\n", ksizes[i]);\r
+\r
+ Mat cpudst;\r
+ cv::blur(img, cpudst, Size(ksizes[i], ksizes[i]));\r
+\r
+ GpuMat gpu1(img);\r
+ GpuMat gpudst;\r
+ cv::gpu::blur(gpu1, gpudst, Size(ksizes[i], ksizes[i]));\r
+\r
+ if (CheckNorm(cpudst, gpudst) != CvTS::OK)\r
+ test_res = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ return test_res;\r
+ }\r
+};\r
+\r
+////////////////////////////////////////////////////////////////////////////////\r
+// cvtColor\r
+class CV_GpuCvtColorTest : public CvTest\r
+{\r
+public:\r
+ CV_GpuCvtColorTest() : CvTest("GPU-CvtColor", "cvtColor") {}\r
+ ~CV_GpuCvtColorTest() {};\r
+\r
+protected:\r
+ void run(int);\r
+ \r
+ int CheckNorm(const Mat& m1, const Mat& m2);\r
+};\r
+\r
+\r
+int CV_GpuCvtColorTest::CheckNorm(const Mat& m1, const Mat& m2)\r
+{\r
+ double ret = norm(m1, m2, NORM_INF);\r
+\r
+ if (ret < std::numeric_limits<double>::epsilon())\r
+ {\r
+ return CvTS::OK;\r
+ }\r
+ else\r
+ {\r
+ ts->printf(CvTS::LOG, "\nNorm: %f\n", ret);\r
+ return CvTS::FAIL_GENERIC;\r
+ }\r
+}\r
+\r
+void CV_GpuCvtColorTest::run( int )\r
+{\r
+ //load image\r
+ cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");\r
+\r
+ if (img.empty())\r
+ {\r
+ ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);\r
+ return;\r
+ }\r
+\r
+ int testResult = CvTS::OK;\r
+ cv::Mat cpuRes;\r
+ cv::gpu::GpuMat gpuImg(img), gpuRes;\r
+ try\r
+ {\r
+ //run tests\r
+ int codes[] = { CV_BGR2RGB, /* CV_RGB2YCrCb, CV_YCrCb2RGB,*/ CV_RGB2RGBA, CV_RGBA2BGRA, CV_BGRA2GRAY, CV_GRAY2RGB, CV_RGB2BGR555/*, CV_BGR5552BGR/*, CV_BGR2BGR565, CV_BGR5652RGB*/};\r
+ const char* codes_str[] = {"CV_BGR2RGB", /*"CV_RGB2YCrCb", "CV_YCrCb2RGB",*/ "CV_RGB2RGBA", "CV_RGBA2BGRA", "CV_BGRA2GRAY", "CV_GRAY2RGB", "CV_RGB2BGR555"/*, "CV_BGR5552BGR"/*, "CV_BGR2BGR565", "CV_BGR5652RGB"*/};\r
+ int codes_num = sizeof(codes) / sizeof(int);\r
+\r
+ for (int i = 0; i < codes_num; ++i)\r
+ {\r
+ ts->printf(CvTS::LOG, "\n%s\n", codes_str[i]);\r
+\r
+ cv::cvtColor(img, cpuRes, codes[i]);\r
+ cv::gpu::cvtColor(gpuImg, gpuRes, codes[i]);\r
+\r
+ if (CheckNorm(cpuRes, gpuRes) == CvTS::OK)\r
+ ts->printf(CvTS::LOG, "\nSUCCESS\n");\r
+ else\r
+ {\r
+ ts->printf(CvTS::LOG, "\nFAIL\n");\r
+ testResult = CvTS::FAIL_GENERIC;\r
+ }\r
+\r
+ img = cpuRes;\r
+ gpuImg = gpuRes;\r
+ }\r
+ }\r
+ catch(const cv::Exception& e)\r
+ {\r
+ if (!check_and_treat_gpu_exception(e, ts))\r
+ throw;\r
+ return;\r
+ }\r
+\r
+ ts->set_failed_test_info(testResult);\r
+}\r
+\r
+/////////////////////////////////////////////////////////////////////////////\r
+/////////////////// tests registration /////////////////////////////////////\r
+/////////////////////////////////////////////////////////////////////////////\r
+\r
+// If we comment some tests, we may foget/miss to uncomment it after.\r
+// Placing all test definitions in one place \r
+// makes us know about what tests are commented.\r
+\r
CV_GpuNppImageThresholdTest CV_GpuNppImageThreshold_test;\r
CV_GpuNppImageResizeTest CV_GpuNppImageResize_test;\r
CV_GpuNppImageCopyMakeBorderTest CV_GpuNppImageCopyMakeBorder_test;\r
projects / profile / ivi / opencv.git / commitdiff