diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index d0105ef409..bfaa6c7c39 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -266,13 +266,13 @@ double cv::gpu::norm(const GpuMat& src1, const GpuMat& src2, int normType) sz.height = src1.rows; int funcIdx = normType >> 1; - Scalar retVal; + double retVal; nppSafeCall( npp_norm_diff_func[funcIdx](src1.ptr(), src1.step, src2.ptr(), src2.step, - sz, retVal.val) ); + sz, &retVal) ); - return retVal[0]; + return retVal; } //////////////////////////////////////////////////////////////////////// @@ -307,10 +307,7 @@ void cv::gpu::flip(const GpuMat& src, GpuMat& dst, int flipCode) Scalar cv::gpu::sum(const GpuMat& src) { - CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4); - - - + CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4); NppiSize sz; sz.width = src.cols; @@ -324,7 +321,7 @@ Scalar cv::gpu::sum(const GpuMat& src) GpuMat buf(1, bufsz, CV_32S); Scalar res; - nppSafeCall( nppiSum_8u_C1R(src.ptr(), src.step, sz, buf.ptr(), res.val) ); + nppSafeCall( nppiSum_8u_C1R(src.ptr(), src.step, sz, buf.ptr(), res.val) ); return res; } else @@ -336,8 +333,6 @@ Scalar cv::gpu::sum(const GpuMat& src) nppSafeCall( nppiSum_8u_C4R(src.ptr(), src.step, sz, buf.ptr(), res.val) ); return res; } - - } //////////////////////////////////////////////////////////////////////// @@ -371,28 +366,54 @@ void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst) { public: Npp32s pLevels[256]; + const Npp32s* pLevels3[3]; + int nValues3[3]; LevelsInit() - { + { + nValues3[0] = nValues3[1] = nValues3[2] = 256; for (int i = 0; i < 256; ++i) pLevels[i] = i; + pLevels3[0] = pLevels3[1] = pLevels3[2] = pLevels; } }; static LevelsInit lvls; int cn = src.channels(); - CV_Assert(src.type() == CV_8UC1); - CV_Assert(lut.depth() == CV_32SC1 && lut.rows * lut.cols == 256 && lut.isContinuous()); + CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC3); + CV_Assert(lut.depth() == CV_8U && (lut.channels() == 1 || lut.channels() == cn) && lut.rows * lut.cols == 256 && lut.isContinuous()); - dst.create(src.size(), src.type()); + dst.create(src.size(), CV_MAKETYPE(lut.depth(), cn)); NppiSize sz; sz.height = src.rows; sz.width = src.cols; + + Mat nppLut; + lut.convertTo(nppLut, CV_32S); - nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr(), src.step, dst.ptr(), dst.step, sz, - lut.ptr(), lvls.pLevels, 256) ); + if (src.type() == CV_8UC1) + { + nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr(), src.step, dst.ptr(), dst.step, sz, + nppLut.ptr(), lvls.pLevels, 256) ); + } + else + { + Mat nppLut3[3]; + const Npp32s* pValues3[3]; + if (nppLut.channels() == 1) + pValues3[0] = pValues3[1] = pValues3[2] = nppLut.ptr(); + else + { + cv::split(nppLut, nppLut3); + pValues3[0] = nppLut3[0].ptr(); + pValues3[1] = nppLut3[1].ptr(); + pValues3[2] = nppLut3[2].ptr(); + } + nppSafeCall( nppiLUT_Linear_8u_C3R(src.ptr(), src.step, dst.ptr(), dst.step, sz, + pValues3, lvls.pLevels3, lvls.nValues3) ); + } } #endif /* !defined (HAVE_CUDA) */ \ No newline at end of file diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu index 6954d570a3..418518ef12 100644 --- a/modules/gpu/src/cuda/color.cu +++ b/modules/gpu/src/cuda/color.cu @@ -65,7 +65,7 @@ namespace imgproc template<> struct TypeVec { typedef float3 vec_t; }; template<> struct TypeVec { typedef float4 vec_t; }; - template struct ColorChannel {}; + template struct ColorChannel {}; template<> struct ColorChannel { @@ -86,7 +86,17 @@ namespace imgproc typedef float worktype_f; static __device__ float max() { return 1.f; } static __device__ float half() { return 0.5f; } - }; + }; + + template + __device__ void assignAlpha(typename TypeVec::vec_t& vec, T val) + { + } + template + __device__ void assignAlpha(typename TypeVec::vec_t& vec, T val) + { + vec.w = val; + } } //////////////////////////////////////// SwapChannels ///////////////////////////////////// @@ -96,7 +106,7 @@ namespace imgproc __constant__ int ccoeffs[4]; template - __global__ void swapChannels(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols) + __global__ void swapChannels(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols) { typedef typename TypeVec::vec_t vec_t; @@ -121,8 +131,8 @@ namespace imgproc namespace cv { namespace gpu { namespace improc { - template - void swapChannels_caller(const DevMem2D_& src, const DevMem2D_& dst, int cn, const int* coeffs, cudaStream_t stream) + template + void swapChannels_caller(const DevMem2D& src, const DevMem2D& dst, const int* coeffs, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -130,39 +140,38 @@ namespace cv { namespace gpu { namespace improc grid.x = divUp(src.cols, threads.x); grid.y = divUp(src.rows, threads.y); - cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, cn * sizeof(int)) ); + cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, CN * sizeof(int)) ); - switch (cn) - { - case 3: - imgproc::swapChannels<3><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols); - break; - case 4: - imgproc::swapChannels<4><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols); - break; - default: - cv::gpu::error("Unsupported channels count", __FILE__, __LINE__); - break; - } + imgproc::swapChannels<<>>(src.ptr, src.step, + dst.ptr, dst.step, src.rows, src.cols); if (stream == 0) cudaSafeCall( cudaThreadSynchronize() ); } - void swapChannels_gpu(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream) + void swapChannels_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream) { - swapChannels_caller(src, dst, cn, coeffs, stream); + typedef void (*swapChannels_caller_t)(const DevMem2D& src, const DevMem2D& dst, const int* coeffs, cudaStream_t stream); + static const swapChannels_caller_t swapChannels_callers[] = {swapChannels_caller, swapChannels_caller}; + + swapChannels_callers[cn - 3](src, dst, coeffs, stream); } - void swapChannels_gpu(const DevMem2D_& src, const DevMem2D_& dst, int cn, const int* coeffs, cudaStream_t stream) + void swapChannels_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream) { - swapChannels_caller(src, dst, cn, coeffs, stream); + typedef void (*swapChannels_caller_t)(const DevMem2D& src, const DevMem2D& dst, const int* coeffs, cudaStream_t stream); + static const swapChannels_caller_t swapChannels_callers[] = {swapChannels_caller, swapChannels_caller}; + + swapChannels_callers[cn - 3](src, dst, coeffs, stream); } - void swapChannels_gpu(const DevMem2Df& src, const DevMem2Df& dst, int cn, const int* coeffs, cudaStream_t stream) + void swapChannels_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream) { - swapChannels_caller(src, dst, cn, coeffs, stream); - } + typedef void (*swapChannels_caller_t)(const DevMem2D& src, const DevMem2D& dst, const int* coeffs, cudaStream_t stream); + static const swapChannels_caller_t swapChannels_callers[] = {swapChannels_caller, swapChannels_caller}; + + swapChannels_callers[cn - 3](src, dst, coeffs, stream); + } }}} ////////////////// Various 3/4-channel to 3/4-channel RGB transformations ///////////////// @@ -170,7 +179,7 @@ namespace cv { namespace gpu { namespace improc namespace imgproc { template - __global__ void RGB2RGB(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols, int bidx) + __global__ void RGB2RGB(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx) { typedef typename TypeVec::vec_t src_t; typedef typename TypeVec::vec_t dst_t; @@ -186,8 +195,7 @@ namespace imgproc dst.x = ((const T*)(&src))[bidx]; dst.y = src.y; dst.z = ((const T*)(&src))[bidx ^ 2]; - if (DSTCN == 4) - ((T*)(&dst))[3] = ColorChannel::max(); + assignAlpha(dst, ColorChannel::max()); *(dst_t*)(dst_ + y * dst_step + x * DSTCN) = dst; } @@ -196,8 +204,8 @@ namespace imgproc namespace cv { namespace gpu { namespace improc { - template - void RGB2RGB_caller(const DevMem2D_& src, int srccn, const DevMem2D_& dst, int dstcn, int bidx, cudaStream_t stream) + template + void RGB2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -205,171 +213,248 @@ namespace cv { namespace gpu { namespace improc grid.x = divUp(src.cols, threads.x); grid.y = divUp(src.rows, threads.y); - switch (dstcn) - { - case 3: - switch (srccn) - { - case 3: - { - int coeffs[] = {2, 1, 0}; - cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, 3 * sizeof(int)) ); - imgproc::swapChannels<3><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols); - break; - } - case 4: - imgproc::RGB2RGB<4, 3><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), - src.rows, src.cols, bidx); - break; - default: - cv::gpu::error("Unsupported channels count", __FILE__, __LINE__); - break; - } - break; - case 4: - switch (srccn) - { - case 3: - imgproc::RGB2RGB<3, 4><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), - src.rows, src.cols, bidx); - break; - case 4: - { - int coeffs[] = {2, 1, 0, 3}; - cudaSafeCall( cudaMemcpyToSymbol(imgproc::ccoeffs, coeffs, 4 * sizeof(int)) ); - imgproc::swapChannels<4><<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols); - break; - } - default: - cv::gpu::error("Unsupported channels count", __FILE__, __LINE__); - break; - } - break; - default: - cv::gpu::error("Unsupported channels count", __FILE__, __LINE__); - break; - } + imgproc::RGB2RGB<<>>(src.ptr, src.step, + dst.ptr, dst.step, src.rows, src.cols, bidx); if (stream == 0) cudaSafeCall( cudaThreadSynchronize() ); } - void RGB2RGB_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream) + void RGB2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream) { - RGB2RGB_caller(src, srccn, dst, dstcn, bidx, stream); + typedef void (*RGB2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); + static const RGB2RGB_caller_t RGB2RGB_callers[2][2] = + { + {RGB2RGB_caller, RGB2RGB_caller}, + {RGB2RGB_caller, RGB2RGB_caller} + }; + + RGB2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream); } - void RGB2RGB_gpu(const DevMem2D_& src, int srccn, const DevMem2D_& dst, int dstcn, int bidx, cudaStream_t stream) + void RGB2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream) { - RGB2RGB_caller(src, srccn, dst, dstcn, bidx, stream); + typedef void (*RGB2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); + static const RGB2RGB_caller_t RGB2RGB_callers[2][2] = + { + {RGB2RGB_caller, RGB2RGB_caller}, + {RGB2RGB_caller, RGB2RGB_caller} + }; + + RGB2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream); } - void RGB2RGB_gpu(const DevMem2Df& src, int srccn, const DevMem2Df& dst, int dstcn, int bidx, cudaStream_t stream) + void RGB2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream) { - RGB2RGB_caller(src, srccn, dst, dstcn, bidx, stream); + typedef void (*RGB2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); + static const RGB2RGB_caller_t RGB2RGB_callers[2][2] = + { + {RGB2RGB_caller, RGB2RGB_caller}, + {RGB2RGB_caller, RGB2RGB_caller} + }; + + RGB2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream); } }}} /////////// Transforming 16-bit (565 or 555) RGB to/from 24/32-bit (888[8]) RGB ////////// -//namespace imgproc -//{ -// struct RGB5x52RGB -// { -// typedef uchar channel_type; -// -// RGB5x52RGB(int _dstcn, int _blueIdx, int _greenBits) -// : dstcn(_dstcn), blueIdx(_blueIdx), greenBits(_greenBits) {} -// -// void operator()(const uchar* src, uchar* dst, int n) const -// { -// int dcn = dstcn, bidx = blueIdx; -// if( greenBits == 6 ) -// for( int i = 0; i < n; i++, dst += dcn ) -// { -// unsigned t = ((const unsigned short*)src)[i]; -// dst[bidx] = (uchar)(t << 3); -// dst[1] = (uchar)((t >> 3) & ~3); -// dst[bidx ^ 2] = (uchar)((t >> 8) & ~7); -// if( dcn == 4 ) -// dst[3] = 255; -// } -// else -// for( int i = 0; i < n; i++, dst += dcn ) -// { -// unsigned t = ((const unsigned short*)src)[i]; -// dst[bidx] = (uchar)(t << 3); -// dst[1] = (uchar)((t >> 2) & ~7); -// dst[bidx ^ 2] = (uchar)((t >> 7) & ~7); -// if( dcn == 4 ) -// dst[3] = t & 0x8000 ? 255 : 0; -// } -// } -// -// int dstcn, blueIdx, greenBits; -// }; -// -// -// struct RGB2RGB5x5 -// { -// typedef uchar channel_type; -// -// RGB2RGB5x5(int _srccn, int _blueIdx, int _greenBits) -// : srccn(_srccn), blueIdx(_blueIdx), greenBits(_greenBits) {} -// -// void operator()(const uchar* src, uchar* dst, int n) const -// { -// int scn = srccn, bidx = blueIdx; -// if( greenBits == 6 ) -// for( int i = 0; i < n; i++, src += scn ) -// { -// ((unsigned short*)dst)[i] = (unsigned short)((src[bidx] >> 3)|((src[1]&~3) << 3)|((src[bidx^2]&~7) << 8)); -// } -// else if( scn == 3 ) -// for( int i = 0; i < n; i++, src += 3 ) -// { -// ((unsigned short*)dst)[i] = (unsigned short)((src[bidx] >> 3)|((src[1]&~7) << 2)|((src[bidx^2]&~7) << 7)); -// } -// else -// for( int i = 0; i < n; i++, src += 4 ) -// { -// ((unsigned short*)dst)[i] = (unsigned short)((src[bidx] >> 3)|((src[1]&~7) << 2)| -// ((src[bidx^2]&~7) << 7)|(src[3] ? 0x8000 : 0)); -// } -// } -// -// int srccn, blueIdx, greenBits; -// }; -//} -// -//namespace cv { namespace gpu { namespace impl -//{ -//}}} - -///////////////////////////////// Grayscale to Color //////////////////////////////// - namespace imgproc { - template - __global__ void Gray2RGB_3(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols) + template struct RGB5x52RGBConverter {}; + + template struct RGB5x52RGBConverter<5, DSTCN> { + typedef typename TypeVec::vec_t dst_t; + + static __device__ dst_t cvt(unsigned int src, int bidx) + { + dst_t dst; + + ((uchar*)(&dst))[bidx] = (uchar)(src << 3); + dst.y = (uchar)((src >> 2) & ~7); + ((uchar*)(&dst))[bidx ^ 2] = (uchar)((src >> 7) & ~7); + assignAlpha(dst, (uchar)(src & 0x8000 ? 255 : 0)); + + return dst; + } + }; + template struct RGB5x52RGBConverter<6, DSTCN> + { + typedef typename TypeVec::vec_t dst_t; + + static __device__ dst_t cvt(unsigned int src, int bidx) + { + dst_t dst; + + ((uchar*)(&dst))[bidx] = (uchar)(src << 3); + dst.y = (uchar)((src >> 3) & ~3); + ((uchar*)(&dst))[bidx ^ 2] = (uchar)((src >> 8) & ~7); + assignAlpha(dst, (uchar)(255)); + + return dst; + } + }; + + template + __global__ void RGB5x52RGB(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx) + { + typedef typename TypeVec::vec_t dst_t; + const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; if (y < rows && x < cols) { - T src = src_[y * src_step + x]; - T* dst = dst_ + y * dst_step + x * 3; - dst[0] = src; - dst[1] = src; - dst[2] = src; + unsigned int src = *(const unsigned short*)(src_ + y * src_step + (x << 1)); + + *(dst_t*)(dst_ + y * dst_step + x * DSTCN) = RGB5x52RGBConverter::cvt(src, bidx); } } - template - __global__ void Gray2RGB_4(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols) + /*struct RGB5x52RGB { - typedef typename TypeVec::vec_t vec4_t; + typedef uchar channel_type; + + RGB5x52RGB(int _dstcn, int _blueIdx, int _greenBits) + : dstcn(_dstcn), blueIdx(_blueIdx), greenBits(_greenBits) {} + + void operator()(const uchar* src, uchar* dst, int n) const + { + int dcn = dstcn, bidx = blueIdx; + if( greenBits == 6 ) + for( int i = 0; i < n; i++, dst += dcn ) + { + unsigned t = ((const unsigned short*)src)[i]; + dst[bidx] = (uchar)(t << 3); + dst[1] = (uchar)((t >> 3) & ~3); + dst[bidx ^ 2] = (uchar)((t >> 8) & ~7); + if( dcn == 4 ) + dst[3] = 255; + } + else + for( int i = 0; i < n; i++, dst += dcn ) + { + unsigned t = ((const unsigned short*)src)[i]; + dst[bidx] = (uchar)(t << 3); + dst[1] = (uchar)((t >> 2) & ~7); + dst[bidx ^ 2] = (uchar)((t >> 7) & ~7); + if( dcn == 4 ) + dst[3] = t & 0x8000 ? 255 : 0; + } + } + + int dstcn, blueIdx, greenBits; + };*/ + + template struct RGB2RGB5x5Converter {}; + + template struct RGB2RGB5x5Converter + { + static __device__ unsigned short cvt(const uchar* src_ptr, int bidx) + { + return (unsigned short)((src_ptr[bidx] >> 3) | ((src_ptr[1] & ~3) << 3) | ((src_ptr[bidx^2] & ~7) << 8)); + } + }; + template<> struct RGB2RGB5x5Converter<3, 5> + { + static __device__ unsigned short cvt(const uchar* src_ptr, int bidx) + { + return (unsigned short)((src_ptr[bidx] >> 3) | ((src_ptr[1] & ~7) << 2) | ((src_ptr[bidx^2] & ~7) << 7)); + } + }; + template<> struct RGB2RGB5x5Converter<4, 5> + { + static __device__ unsigned short cvt(const uchar* src_ptr, int bidx) + { + return (unsigned short)((src_ptr[bidx] >> 3) | ((src_ptr[1] & ~7) << 2) | ((src_ptr[bidx^2] & ~7) << 7)|(src_ptr[3] ? 0x8000 : 0)); + } + }; + + template + __global__ void RGB2RGB5x5(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx) + { + typedef typename TypeVec::vec_t src_t; + + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (y < rows && x < cols) + { + src_t src = *(src_t*)(src_ + y * src_step + x * SRCCN); + + *(unsigned short*)(dst_ + y * dst_step + (x << 1)) = RGB2RGB5x5Converter::cvt((const uchar*)(&src), bidx); + } + } +} + +namespace cv { namespace gpu { namespace improc +{ + template + void RGB5x52RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream) + { + dim3 threads(32, 8, 1); + dim3 grid(1, 1, 1); + + grid.x = divUp(src.cols, threads.x); + grid.y = divUp(src.rows, threads.y); + + imgproc::RGB5x52RGB<<>>(src.ptr, src.step, + dst.ptr, dst.step, src.rows, src.cols, bidx); + + if (stream == 0) + cudaSafeCall( cudaThreadSynchronize() ); + } + + void RGB5x52RGB_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream) + { + typedef void (*RGB5x52RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); + static const RGB5x52RGB_caller_t RGB5x52RGB_callers[2][2] = + { + {RGB5x52RGB_caller<5, 3>, RGB5x52RGB_caller<5, 4>}, + {RGB5x52RGB_caller<6, 3>, RGB5x52RGB_caller<6, 4>} + }; + + RGB5x52RGB_callers[green_bits - 5][dstcn - 5](src, dst, bidx, stream); + } + + template + void RGB2RGB5x5_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream) + { + dim3 threads(32, 8, 1); + dim3 grid(1, 1, 1); + + grid.x = divUp(src.cols, threads.x); + grid.y = divUp(src.rows, threads.y); + + imgproc::RGB2RGB5x5<<>>(src.ptr, src.step, + dst.ptr, dst.step, src.rows, src.cols, bidx); + + if (stream == 0) + cudaSafeCall( cudaThreadSynchronize() ); + } + + void RGB2RGB5x5_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int green_bits, int bidx, cudaStream_t stream) + { + typedef void (*RGB2RGB5x5_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream); + static const RGB2RGB5x5_caller_t RGB2RGB5x5_callers[2][2] = + { + {RGB2RGB5x5_caller<3, 5>, RGB2RGB5x5_caller<3, 6>}, + {RGB2RGB5x5_caller<4, 5>, RGB2RGB5x5_caller<4, 6>} + }; + + RGB2RGB5x5_callers[srccn - 3][green_bits - 5](src, dst, bidx, stream); + } +}}} + +///////////////////////////////// Grayscale to Color //////////////////////////////// + +namespace imgproc +{ + template + __global__ void Gray2RGB(const T* src_, size_t src_step, T* dst_, size_t dst_step, int rows, int cols) + { + typedef typename TypeVec::vec_t dst_t; const int x = blockDim.x * blockIdx.x + threadIdx.x; const int y = blockDim.y * blockIdx.y + threadIdx.y; @@ -377,12 +462,12 @@ namespace imgproc if (y < rows && x < cols) { T src = src_[y * src_step + x]; - vec4_t dst; + dst_t dst; dst.x = src; dst.y = src; dst.z = src; - dst.w = ColorChannel::max(); - *(vec4_t*)(dst_ + y * dst_step + (x << 2)) = dst; + assignAlpha(dst, ColorChannel::max()); + *(dst_t*)(dst_ + y * dst_step + x * DSTCN) = dst; } } @@ -412,8 +497,8 @@ namespace imgproc namespace cv { namespace gpu { namespace improc { - template - void Gray2RGB_caller(const DevMem2D_& src, const DevMem2D_& dst, int dstcn, cudaStream_t stream) + template + void Gray2RGB_caller(const DevMem2D_& src, const DevMem2D_& dst, cudaStream_t stream) { dim3 threads(32, 8, 1); dim3 grid(1, 1, 1); @@ -421,18 +506,8 @@ namespace cv { namespace gpu { namespace improc grid.x = divUp(src.cols, threads.x); grid.y = divUp(src.rows, threads.y); - switch (dstcn) - { - case 3: - imgproc::Gray2RGB_3<<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols); - break; - case 4: - imgproc::Gray2RGB_4<<>>(src.ptr, src.step / sizeof(T), dst.ptr, dst.step / sizeof(T), src.rows, src.cols); - break; - default: - cv::gpu::error("Unsupported channels count", __FILE__, __LINE__); - break; - } + imgproc::Gray2RGB<<>>(src.ptr, src.step / sizeof(T), + dst.ptr, dst.step / sizeof(T), src.rows, src.cols); if (stream == 0) cudaSafeCall( cudaThreadSynchronize() ); @@ -440,17 +515,26 @@ namespace cv { namespace gpu { namespace improc void Gray2RGB_gpu(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream) { - Gray2RGB_caller(src, dst, dstcn, stream); + typedef void (*Gray2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream); + static const Gray2RGB_caller_t Gray2RGB_callers[] = {Gray2RGB_caller, Gray2RGB_caller}; + + Gray2RGB_callers[dstcn - 3](src, dst, stream); } void Gray2RGB_gpu(const DevMem2D_& src, const DevMem2D_& dst, int dstcn, cudaStream_t stream) { - Gray2RGB_caller(src, dst, dstcn, stream); + typedef void (*Gray2RGB_caller_t)(const DevMem2D_& src, const DevMem2D_& dst, cudaStream_t stream); + static const Gray2RGB_caller_t Gray2RGB_callers[] = {Gray2RGB_caller, Gray2RGB_caller}; + + Gray2RGB_callers[dstcn - 3](src, dst, stream); } void Gray2RGB_gpu(const DevMem2Df& src, const DevMem2Df& dst, int dstcn, cudaStream_t stream) { - Gray2RGB_caller(src, dst, dstcn, stream); + typedef void (*Gray2RGB_caller_t)(const DevMem2Df& src, const DevMem2Df& dst, cudaStream_t stream); + static const Gray2RGB_caller_t Gray2RGB_callers[] = {Gray2RGB_caller, Gray2RGB_caller}; + + Gray2RGB_callers[dstcn - 3](src, dst, stream); } }}} diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp index 8a207f2955..85766e9712 100644 --- a/modules/gpu/src/imgproc_gpu.cpp +++ b/modules/gpu/src/imgproc_gpu.cpp @@ -81,13 +81,16 @@ namespace cv { namespace gpu void reprojectImageTo3D_gpu(const DevMem2D& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream); void reprojectImageTo3D_gpu(const DevMem2D_& disp, const DevMem2Df& xyzw, const float* q, const cudaStream_t& stream); - void swapChannels_gpu(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream); - void swapChannels_gpu(const DevMem2D_& src, const DevMem2D_& dst, int cn, const int* coeffs, cudaStream_t stream); - void swapChannels_gpu(const DevMem2Df& src, const DevMem2Df& dst, int cn, const int* coeffs, cudaStream_t stream); + void swapChannels_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream); + void swapChannels_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream); + void swapChannels_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int cn, const int* coeffs, cudaStream_t stream); - void RGB2RGB_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream); - void RGB2RGB_gpu(const DevMem2D_& src, int srccn, const DevMem2D_& dst, int dstcn, int bidx, cudaStream_t stream); - void RGB2RGB_gpu(const DevMem2Df& src, int srccn, const DevMem2Df& dst, int dstcn, int bidx, cudaStream_t stream); + void RGB2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream); + void RGB2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream); + void RGB2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream); + + void RGB5x52RGB_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, int dstcn, int bidx, cudaStream_t stream); + void RGB2RGB5x5_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int green_bits, int bidx, cudaStream_t stream); void Gray2RGB_gpu(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream); void Gray2RGB_gpu(const DevMem2D_& src, const DevMem2D_& dst, int dstcn, cudaStream_t stream); @@ -245,38 +248,36 @@ namespace out.create(sz, CV_MAKETYPE(depth, dcn)); if( depth == CV_8U ) - improc::RGB2RGB_gpu((DevMem2D)src, scn, (DevMem2D)out, dcn, bidx, stream); + improc::RGB2RGB_gpu_8u(src, scn, out, dcn, bidx, stream); else if( depth == CV_16U ) - improc::RGB2RGB_gpu((DevMem2D_)src, scn, (DevMem2D_)out, dcn, bidx, stream); + improc::RGB2RGB_gpu_16u(src, scn, out, dcn, bidx, stream); else - improc::RGB2RGB_gpu((DevMem2Df)src, scn, (DevMem2Df)out, dcn, bidx, stream); + improc::RGB2RGB_gpu_32f(src, scn, out, dcn, bidx, stream); break; - //case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555: - //case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555: - // CV_Assert( (scn == 3 || scn == 4) && depth == CV_8U ); - // dst.create(sz, CV_8UC2); - // - // CvtColorLoop(src, dst, RGB2RGB5x5(scn, - // code == CV_BGR2BGR565 || code == CV_BGR2BGR555 || - // code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2, - // code == CV_BGR2BGR565 || code == CV_RGB2BGR565 || - // code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5 // green bits - // )); - // break; + case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555: + case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555: + CV_Assert( (scn == 3 || scn == 4) && depth == CV_8U ); + out.create(sz, CV_8UC2); + + improc::RGB2RGB5x5_gpu(src, scn, out, code == CV_BGR2BGR565 || code == CV_RGB2BGR565 || + code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5, + code == CV_BGR2BGR565 || code == CV_BGR2BGR555 || + code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2, + stream); + break; //case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: //case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: // if(dcn <= 0) dcn = 3; // CV_Assert( (dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U ); - // dst.create(sz, CV_MAKETYPE(depth, dcn)); - // - // CvtColorLoop(src, dst, RGB5x52RGB(dcn, - // code == CV_BGR5652BGR || code == CV_BGR5552BGR || - // code == CV_BGR5652BGRA || code == CV_BGR5552BGRA ? 0 : 2, // blue idx - // code == CV_BGR5652BGR || code == CV_BGR5652RGB || - // code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5 // green bits - // )); + // out.create(sz, CV_MAKETYPE(depth, dcn)); + + // improc::RGB5x52RGB_gpu(src, code == CV_BGR2BGR565 || code == CV_RGB2BGR565 || + // code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5, out, dcn, + // code == CV_BGR2BGR565 || code == CV_BGR2BGR555 || + // code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2, + // stream); // break; case CV_BGR2GRAY: case CV_BGRA2GRAY: case CV_RGB2GRAY: case CV_RGBA2GRAY: @@ -329,7 +330,7 @@ namespace nppSafeCall( nppiRGBToYCbCr_8u_C3R(src.ptr(), src.step, out.ptr(), out.step, nppsz) ); { static int coeffs[] = {0, 2, 1}; - improc::swapChannels_gpu((DevMem2D)out, (DevMem2D)out, 3, coeffs, 0); + improc::swapChannels_gpu_8u(out, out, 3, coeffs, 0); } break; @@ -341,7 +342,7 @@ namespace { static int coeffs[] = {0, 2, 1}; GpuMat src1(src.size(), src.type()); - improc::swapChannels_gpu((DevMem2D)src, (DevMem2D)src1, 3, coeffs, 0); + improc::swapChannels_gpu_8u(src, src1, 3, coeffs, 0); nppSafeCall( nppiYCbCrToRGB_8u_C3R(src1.ptr(), src1.step, out.ptr(), out.step, nppsz) ); } break; diff --git a/tests/gpu/src/arithm.cpp b/tests/gpu/src/arithm.cpp index b9033f0370..f4168436a3 100644 --- a/tests/gpu/src/arithm.cpp +++ b/tests/gpu/src/arithm.cpp @@ -1,585 +1,545 @@ -/*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 -#include -#include -#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::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::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/////////////////////////////////////////////////////////////////////////////////////// +// +// 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 +#include +#include +#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::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::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; + } + } + } + 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 && mat1.type() != CV_8UC3) + { + ts->printf(CvTS::LOG, "\nUnsupported type\n"); + return CvTS::OK; + } + + cv::Mat lut(1, 256, CV_8UC1); + 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); + + 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. + CV_GpuNppImageAddTest CV_GpuNppImageAdd_test; CV_GpuNppImageSubtractTest CV_GpuNppImageSubtract_test; CV_GpuNppImageMultiplyTest CV_GpuNppImageMultiply_test; diff --git a/tests/gpu/src/gputest_main.cpp b/tests/gpu/src/gputest_main.cpp index 2a2be4d607..143d380fcc 100644 --- a/tests/gpu/src/gputest_main.cpp +++ b/tests/gpu/src/gputest_main.cpp @@ -46,6 +46,18 @@ CvTS test_system; 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 }; diff --git a/tests/gpu/src/imgproc_gpu.cpp b/tests/gpu/src/imgproc_gpu.cpp index 19e499329e..0f8a6aed2e 100644 --- a/tests/gpu/src/imgproc_gpu.cpp +++ b/tests/gpu/src/imgproc_gpu.cpp @@ -1,547 +1,544 @@ -/*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 -#include -#include -#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::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::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/////////////////////////////////////////////////////////////////////////////////////// +// +// 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 +#include +#include +#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::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])); + + 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-CvtColor", "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::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, CV_RGB2BGR555/*, CV_BGR5552BGR/*, CV_BGR2BGR565, CV_BGR5652RGB*/}; + 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"*/}; + 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. + CV_GpuNppImageThresholdTest CV_GpuNppImageThreshold_test; CV_GpuNppImageResizeTest CV_GpuNppImageResize_test; CV_GpuNppImageCopyMakeBorderTest CV_GpuNppImageCopyMakeBorder_test;