diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index 76e2621860..3ef489813b 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -91,6 +91,12 @@ void cv::gpu::gemm(const GpuMat& src1, const GpuMat& src2, double alpha, const G bool tr2 = (flags & GEMM_2_T) != 0; bool tr3 = (flags & GEMM_3_T) != 0; + if (src1.type() == CV_64FC2) + { + if (tr1 || tr2 || tr3) + CV_Error(CV_StsNotImplemented, "transpose operation doesn't implemented for CV_64FC2 type"); + } + Size src1Size = tr1 ? Size(src1.rows, src1.cols) : src1.size(); Size src2Size = tr2 ? Size(src2.rows, src2.cols) : src2.size(); Size src3Size = tr3 ? Size(src3.rows, src3.cols) : src3.size(); @@ -99,7 +105,7 @@ void cv::gpu::gemm(const GpuMat& src1, const GpuMat& src2, double alpha, const G CV_Assert(src1Size.width == src2Size.height); CV_Assert(src3.empty() || src3Size == dstSize); - dst.create(dstSize, CV_32FC1); + dst.create(dstSize, src1.type()); if (beta != 0) { @@ -149,7 +155,7 @@ void cv::gpu::gemm(const GpuMat& src1, const GpuMat& src2, double alpha, const G { case CV_32FC1: cublasSafeCall( cublasSgemm_v2(handle, transa, transb, tr2 ? src2.rows : src2.cols, tr1 ? src1.cols : src1.rows, tr2 ? src2.cols : src2.rows, - &alphaf, + &alphaf, src2.ptr(), static_cast(src2.step / sizeof(float)), src1.ptr(), static_cast(src1.step / sizeof(float)), &betaf, @@ -158,7 +164,7 @@ void cv::gpu::gemm(const GpuMat& src1, const GpuMat& src2, double alpha, const G case CV_64FC1: cublasSafeCall( cublasDgemm_v2(handle, transa, transb, tr2 ? src2.rows : src2.cols, tr1 ? src1.cols : src1.rows, tr2 ? src2.cols : src2.rows, - &alpha, + &alpha, src2.ptr(), static_cast(src2.step / sizeof(double)), src1.ptr(), static_cast(src1.step / sizeof(double)), &beta, @@ -167,7 +173,7 @@ void cv::gpu::gemm(const GpuMat& src1, const GpuMat& src2, double alpha, const G case CV_32FC2: cublasSafeCall( cublasCgemm_v2(handle, transa, transb, tr2 ? src2.rows : src2.cols, tr1 ? src1.cols : src1.rows, tr2 ? src2.cols : src2.rows, - &alphacf, + &alphacf, src2.ptr(), static_cast(src2.step / sizeof(cuComplex)), src1.ptr(), static_cast(src1.step / sizeof(cuComplex)), &betacf, @@ -176,7 +182,7 @@ void cv::gpu::gemm(const GpuMat& src1, const GpuMat& src2, double alpha, const G case CV_64FC2: cublasSafeCall( cublasZgemm_v2(handle, transa, transb, tr2 ? src2.rows : src2.cols, tr1 ? src1.cols : src1.rows, tr2 ? src2.cols : src2.rows, - &alphac, + &alphac, src2.ptr(), static_cast(src2.step / sizeof(cuDoubleComplex)), src1.ptr(), static_cast(src1.step / sizeof(cuDoubleComplex)), &betac, @@ -208,8 +214,8 @@ void cv::gpu::transpose(const GpuMat& src, GpuMat& dst, Stream& s) sz.width = src.cols; sz.height = src.rows; - nppSafeCall( nppiTranspose_8u_C1R(src.ptr(), static_cast(src.step), - dst.ptr(), static_cast(dst.step), sz) ); + nppSafeCall( nppiTranspose_8u_C1R(src.ptr(), static_cast(src.step), + dst.ptr(), static_cast(dst.step), sz) ); } else if (src.elemSize() == 4) { @@ -219,7 +225,7 @@ void cv::gpu::transpose(const GpuMat& src, GpuMat& dst, Stream& s) sz.width = src.cols; sz.height = src.rows; - ncvSafeCall( nppiStTranspose_32u_C1R(const_cast(src.ptr()), static_cast(src.step), + ncvSafeCall( nppiStTranspose_32u_C1R(const_cast(src.ptr()), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz) ); } else // if (src.elemSize() == 8) @@ -230,8 +236,8 @@ void cv::gpu::transpose(const GpuMat& src, GpuMat& dst, Stream& s) sz.width = src.cols; sz.height = src.rows; - ncvSafeCall( nppiStTranspose_64u_C1R(const_cast(src.ptr()), static_cast(src.step), - dst.ptr(), static_cast(dst.step), sz) ); + ncvSafeCall( nppiStTranspose_64u_C1R(const_cast(src.ptr()), static_cast(src.step), + dst.ptr(), static_cast(dst.step), sz) ); } if (stream == 0) @@ -285,7 +291,7 @@ void cv::gpu::flip(const GpuMat& src, GpuMat& dst, int flipCode, Stream& stream) { typedef void (*func_t)(const GpuMat& src, GpuMat& dst, int flipCode, cudaStream_t stream); - static const func_t funcs[6][4] = + static const func_t funcs[6][4] = { {NppMirror::call, 0, NppMirror::call, NppMirror::call}, {0,0,0,0}, @@ -345,7 +351,7 @@ void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst, Stream& s) if (src.type() == CV_8UC1) { - nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr(), static_cast(src.step), + nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz, nppLut.ptr(), lvls.pLevels, 256) ); } else @@ -361,7 +367,7 @@ void cv::gpu::LUT(const GpuMat& src, const Mat& lut, GpuMat& dst, Stream& s) pValues3[1] = nppLut3[1].ptr(); pValues3[2] = nppLut3[2].ptr(); } - nppSafeCall( nppiLUT_Linear_8u_C3R(src.ptr(), static_cast(src.step), + nppSafeCall( nppiLUT_Linear_8u_C3R(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), sz, pValues3, lvls.pLevels3, lvls.nValues3) ); } @@ -408,9 +414,9 @@ void cv::gpu::magnitudeSqr(const GpuMat& src, GpuMat& dst, Stream& stream) //////////////////////////////////////////////////////////////////////// // Polar <-> Cart -namespace cv { namespace gpu { namespace device +namespace cv { namespace gpu { namespace device { - namespace mathfunc + namespace mathfunc { void cartToPolar_gpu(DevMem2Df x, DevMem2Df y, DevMem2Df mag, bool magSqr, DevMem2Df angle, bool angleInDegrees, cudaStream_t stream); void polarToCart_gpu(DevMem2Df mag, DevMem2Df angle, DevMem2Df x, DevMem2Df y, bool angleInDegrees, cudaStream_t stream); diff --git a/modules/gpu/src/cuda/element_operations.cu b/modules/gpu/src/cuda/element_operations.cu index f5e8459091..4f4efc27dd 100644 --- a/modules/gpu/src/cuda/element_operations.cu +++ b/modules/gpu/src/cuda/element_operations.cu @@ -1672,40 +1672,53 @@ namespace cv { namespace gpu { namespace device template::is_signed> struct PowOp : unary_function { - float power; - PowOp(float power_) : power(power_) {} + const float power; - __device__ __forceinline__ T operator()(const T& e) const + PowOp(double power_) : power(static_cast(power_)) {} + + __device__ __forceinline__ T operator()(T e) const { return saturate_cast(__powf((float)e, power)); } }; - template struct PowOp : unary_function { - float power; - PowOp(float power_) : power(power_) {} + const float power; + + PowOp(double power_) : power(static_cast(power_)) {} - __device__ __forceinline__ float operator()(const T& e) const + __device__ __forceinline__ T operator()(T e) const { T res = saturate_cast(__powf((float)e, power)); - if ( (e < 0) && (1 & (int)power) ) - res *= -1; + if ((e < 0) && (1 & static_cast(power))) + res *= -1; + return res; } }; - template<> struct PowOp : unary_function { - float power; - PowOp(float power_) : power(power_) {} + const float power; + + PowOp(double power_) : power(static_cast(power_)) {} - __device__ __forceinline__ float operator()(const float& e) const + __device__ __forceinline__ float operator()(float e) const { return __powf(::fabs(e), power); } }; + template<> struct PowOp : unary_function + { + const double power; + + PowOp(double power_) : power(power_) {} + + __device__ __forceinline__ double operator()(double e) const + { + return ::pow(::fabs(e), power); + } + }; namespace detail { @@ -1733,17 +1746,18 @@ namespace cv { namespace gpu { namespace device }; template - void pow_caller(const DevMem2Db& src, float power, DevMem2Db dst, cudaStream_t stream) + void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream) { cv::gpu::device::transform((DevMem2D_)src, (DevMem2D_)dst, PowOp(power), WithOutMask(), stream); } - template void pow_caller(const DevMem2Db& src, float power, DevMem2Db dst, cudaStream_t stream); - template void pow_caller(const DevMem2Db& src, float power, DevMem2Db dst, cudaStream_t stream); - template void pow_caller(const DevMem2Db& src, float power, DevMem2Db dst, cudaStream_t stream); - template void pow_caller(const DevMem2Db& src, float power, DevMem2Db dst, cudaStream_t stream); - template void pow_caller(const DevMem2Db& src, float power, DevMem2Db dst, cudaStream_t stream); - template void pow_caller(const DevMem2Db& src, float power, DevMem2Db dst, cudaStream_t stream); + template void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream); + template void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream); + template void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream); + template void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream); + template void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream); + template void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream); + template void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream); ////////////////////////////////////////////////////////////////////////// // addWeighted diff --git a/modules/gpu/src/element_operations.cpp b/modules/gpu/src/element_operations.cpp index 9abe173ae8..94eefe90cb 100644 --- a/modules/gpu/src/element_operations.cpp +++ b/modules/gpu/src/element_operations.cpp @@ -1301,50 +1301,26 @@ void cv::gpu::compare(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, int c }; CV_Assert(src1.size() == src2.size() && src1.type() == src2.type()); + CV_Assert(cmpop >= CMP_EQ && cmpop <= CMP_NE); - int code; - const GpuMat* psrc1; - const GpuMat* psrc2; - - switch (cmpop) - { - case CMP_EQ: - code = 0; - psrc1 = &src1; - psrc2 = &src2; - break; - case CMP_GE: - code = 3; - psrc1 = &src2; - psrc2 = &src1; - break; - case CMP_GT: - code = 2; - psrc1 = &src2; - psrc2 = &src1; - break; - case CMP_LE: - code = 3; - psrc1 = &src1; - psrc2 = &src2; - break; - case CMP_LT: - code = 2; - psrc1 = &src1; - psrc2 = &src2; - break; - case CMP_NE: - code = 1; - psrc1 = &src1; - psrc2 = &src2; - break; - default: - CV_Error(CV_StsBadFlag, "Incorrect compare operation"); + static const int codes[] = + { + 0, 2, 3, 2, 3, 1 + }; + + const GpuMat* psrc1[] = + { + &src1, &src2, &src2, &src1, &src1, &src1 + }; + + const GpuMat* psrc2[] = + { + &src2, &src1, &src1, &src2, &src2, &src2 }; dst.create(src1.size(), CV_MAKE_TYPE(CV_8U, src1.channels())); - funcs[src1.depth()][code](psrc1->reshape(1), psrc2->reshape(1), dst.reshape(1), StreamAccessor::getStream(stream)); + funcs[src1.depth()][codes[cmpop]](psrc1[cmpop]->reshape(1), psrc2[cmpop]->reshape(1), dst.reshape(1), StreamAccessor::getStream(stream)); } @@ -1944,26 +1920,25 @@ double cv::gpu::threshold(const GpuMat& src, GpuMat& dst, double thresh, double namespace cv { namespace gpu { namespace device { template - void pow_caller(const DevMem2Db& src, float power, DevMem2Db dst, cudaStream_t stream); + void pow_caller(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream); }}} void cv::gpu::pow(const GpuMat& src, double power, GpuMat& dst, Stream& stream) { - using namespace ::cv::gpu::device; - - CV_Assert(src.depth() != CV_64F); - dst.create(src.size(), src.type()); + using namespace cv::gpu::device; - typedef void (*caller_t)(const DevMem2Db& src, float power, DevMem2Db dst, cudaStream_t stream); + typedef void (*func_t)(DevMem2Db src, double power, DevMem2Db dst, cudaStream_t stream); - static const caller_t callers[] = + static const func_t funcs[] = { pow_caller, pow_caller, pow_caller, pow_caller, - pow_caller, pow_caller + pow_caller, pow_caller, pow_caller }; - callers[src.depth()](src.reshape(1), (float)power, dst.reshape(1), StreamAccessor::getStream(stream)); + dst.create(src.size(), src.type()); + + funcs[src.depth()](src.reshape(1), power, dst.reshape(1), StreamAccessor::getStream(stream)); } //////////////////////////////////////////////////////////////////////// @@ -2052,27 +2027,11 @@ namespace cv { namespace gpu { namespace device void cv::gpu::addWeighted(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst, int dtype, Stream& stream) { - using namespace ::cv::gpu::device; - - CV_Assert(src1.size() == src2.size()); - CV_Assert(src1.type() == src2.type() || (dtype >= 0 && src1.channels() == src2.channels())); - - dtype = dtype >= 0 ? CV_MAKETYPE(dtype, src1.channels()) : src1.type(); - - dst.create(src1.size(), dtype); - - const GpuMat* psrc1 = &src1; - const GpuMat* psrc2 = &src2; - - if (src1.depth() > src2.depth()) - { - std::swap(psrc1, psrc2); - std::swap(alpha, beta); - } + using namespace cv::gpu::device; - typedef void (*caller_t)(const DevMem2Db& src1, double alpha, const DevMem2Db& src2, double beta, double gamma, const DevMem2Db& dst, cudaStream_t stream); + typedef void (*func_t)(const DevMem2Db& src1, double alpha, const DevMem2Db& src2, double beta, double gamma, const DevMem2Db& dst, cudaStream_t stream); - static const caller_t callers[7][7][7] = + static const func_t funcs[7][7][7] = { { { @@ -2531,7 +2490,26 @@ void cv::gpu::addWeighted(const GpuMat& src1, double alpha, const GpuMat& src2, } }; - callers[psrc1->depth()][psrc2->depth()][dst.depth()](psrc1->reshape(1), alpha, psrc2->reshape(1), beta, gamma, dst.reshape(1), StreamAccessor::getStream(stream)); + CV_Assert(src1.size() == src2.size()); + CV_Assert(src1.type() == src2.type() || (dtype >= 0 && src1.channels() == src2.channels())); + + dtype = dtype >= 0 ? CV_MAKETYPE(dtype, src1.channels()) : src1.type(); + + dst.create(src1.size(), dtype); + + const GpuMat* psrc1 = &src1; + const GpuMat* psrc2 = &src2; + + if (src1.depth() > src2.depth()) + { + std::swap(psrc1, psrc2); + std::swap(alpha, beta); + } + + const func_t func = funcs[psrc1->depth()][psrc2->depth()][dst.depth()]; + CV_Assert(func != 0); + + func(psrc1->reshape(1), alpha, psrc2->reshape(1), beta, gamma, dst.reshape(1), StreamAccessor::getStream(stream)); } #endif diff --git a/modules/gpu/test/main.cpp b/modules/gpu/test/main.cpp index e6f1ba054f..3370fbce1d 100644 --- a/modules/gpu/test/main.cpp +++ b/modules/gpu/test/main.cpp @@ -52,7 +52,7 @@ using namespace cvtest; using namespace testing; void print_info() -{ +{ printf("\n"); #if defined _WIN32 # if defined _WIN64 @@ -78,9 +78,9 @@ void print_info() int driver; cudaDriverGetVersion(&driver); - printf("CUDA Driver version: %d\n", driver); - printf("CUDA Runtime version: %d\n", CUDART_VERSION); - printf("CUDA device count: %d\n\n", deviceCount); + printf("CUDA Driver version: %d\n", driver); + printf("CUDA Runtime version: %d\n", CUDART_VERSION); + printf("CUDA device count: %d\n\n", deviceCount); for (int i = 0; i < deviceCount; ++i) { @@ -96,7 +96,7 @@ void print_info() else puts(" This device is NOT compatible with current GPU module build\n"); } - + puts("GPU module was compiled for the following GPU archs:"); printf(" BIN: %s\n", CUDA_ARCH_BIN); printf(" PTX: %s\n\n", CUDA_ARCH_PTX); diff --git a/modules/gpu/test/test_arithm.cpp b/modules/gpu/test/test_core.cpp similarity index 64% rename from modules/gpu/test/test_arithm.cpp rename to modules/gpu/test/test_core.cpp index 0bf806b2a1..f9de3b98a7 100644 --- a/modules/gpu/test/test_arithm.cpp +++ b/modules/gpu/test/test_core.cpp @@ -41,8 +41,6 @@ #include "precomp.hpp" -#ifdef HAVE_CUDA - //////////////////////////////////////////////////////////////////////////////// // Add_Array @@ -74,12 +72,6 @@ PARAM_TEST_CASE(Add_Array, cv::gpu::DeviceInfo, cv::Size, std::pairget_rng(); - - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); - - mat1 = randomMat(rng, size, type, 5, 16, false); - mat2 = randomMat(rng, size, type, 5, 16, false); - - val = cv::Scalar(rng.uniform(1, 3), rng.uniform(1, 3), rng.uniform(1, 3), rng.uniform(1, 3)); } }; -//////////////////////////////////////////////////////////////////////////////// -// transpose - -struct Transpose : ArithmTestBase {}; - -TEST_P(Transpose, Accuracy) +TEST_P(Pow, Accuracy) { - cv::Mat dst_gold; - cv::transpose(mat1, dst_gold); - - cv::Mat dst; + cv::Mat src = randomMat(size, depth, 0.0, 100.0); + double power = randomDouble(2.0, 4.0); - cv::gpu::GpuMat gpuRes; + if (src.depth() < CV_32F) + power = static_cast(power); - cv::gpu::transpose(loadMat(mat1, useRoi), gpuRes); + cv::gpu::GpuMat dst = createMat(size, depth, useRoi); + cv::gpu::pow(loadMat(src, useRoi), power, dst); - gpuRes.download(dst); + cv::Mat dst_gold; + cv::pow(src, power, dst_gold); - EXPECT_MAT_NEAR(dst_gold, dst, 0.0); + EXPECT_MAT_NEAR(dst_gold, dst, depth < CV_32F ? 0.0 : 1e-6); } -INSTANTIATE_TEST_CASE_P(Arithm, Transpose, Combine( - ALL_DEVICES, - Values(CV_8UC1, CV_8UC4, CV_8SC1, CV_8SC4, CV_16UC2, CV_16SC2, CV_32SC1, CV_32SC2, CV_32FC1, CV_32FC2, CV_64FC1), - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, Pow, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + ALL_DEPTH, + WHOLE_SUBMAT)); -//////////////////////////////////////////////////////////////////////////////// -// meanStdDev +////////////////////////////////////////////////////////////////////////////// +// AddWeighted -PARAM_TEST_CASE(MeanStdDev, cv::gpu::DeviceInfo, UseRoi) +PARAM_TEST_CASE(AddWeighted, cv::gpu::DeviceInfo, cv::Size, MatDepth, MatDepth, MatDepth, UseRoi) { cv::gpu::DeviceInfo devInfo; - bool useRoi; - cv::Size size; - cv::Mat mat; + int depth1; + int depth2; + int dst_depth; + bool useRoi; - cv::Scalar mean_gold; - cv::Scalar stddev_gold; virtual void SetUp() { devInfo = GET_PARAM(0); - useRoi = GET_PARAM(1); + size = GET_PARAM(1); + depth1 = GET_PARAM(2); + depth2 = GET_PARAM(3); + dst_depth = GET_PARAM(4); + useRoi = GET_PARAM(5); cv::gpu::setDevice(devInfo.deviceID()); - - cv::RNG& rng = TS::ptr()->get_rng(); - - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); - - mat = randomMat(rng, size, CV_8UC1, 1, 255, false); - - cv::meanStdDev(mat, mean_gold, stddev_gold); } }; -TEST_P(MeanStdDev, Accuracy) +TEST_P(AddWeighted, Accuracy) { - cv::Scalar mean; - cv::Scalar stddev; + cv::Mat src1 = randomMat(size, depth1); + cv::Mat src2 = randomMat(size, depth2); + double alpha = randomDouble(-10.0, 10.0); + double beta = randomDouble(-10.0, 10.0); + double gamma = randomDouble(-10.0, 10.0); - cv::gpu::meanStdDev(loadMat(mat, useRoi), mean, stddev); + cv::gpu::GpuMat dst = createMat(size, dst_depth, useRoi); + cv::gpu::addWeighted(loadMat(src1, useRoi), alpha, loadMat(src2, useRoi), beta, gamma, dst, dst_depth); - EXPECT_NEAR(mean_gold[0], mean[0], 1e-5); - EXPECT_NEAR(mean_gold[1], mean[1], 1e-5); - EXPECT_NEAR(mean_gold[2], mean[2], 1e-5); - EXPECT_NEAR(mean_gold[3], mean[3], 1e-5); + cv::Mat dst_gold; + cv::addWeighted(src1, alpha, src2, beta, gamma, dst_gold, dst_depth); - EXPECT_NEAR(stddev_gold[0], stddev[0], 1e-5); - EXPECT_NEAR(stddev_gold[1], stddev[1], 1e-5); - EXPECT_NEAR(stddev_gold[2], stddev[2], 1e-5); - EXPECT_NEAR(stddev_gold[3], stddev[3], 1e-5); + EXPECT_MAT_NEAR(dst_gold, dst, dst_depth < CV_32F ? 1.0 : 1e-12); } -INSTANTIATE_TEST_CASE_P(Arithm, MeanStdDev, Combine( - ALL_DEVICES, - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, AddWeighted, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + ALL_DEPTH, + ALL_DEPTH, + ALL_DEPTH, + WHOLE_SUBMAT)); -//////////////////////////////////////////////////////////////////////////////// -// normDiff +////////////////////////////////////////////////////////////////////////////// +// GEMM -PARAM_TEST_CASE(NormDiff, cv::gpu::DeviceInfo, NormCode, UseRoi) +PARAM_TEST_CASE(GEMM, cv::gpu::DeviceInfo, cv::Size, MatType, GemmFlags, UseRoi) { cv::gpu::DeviceInfo devInfo; - int normCode; - bool useRoi; - cv::Size size; - cv::Mat mat1, mat2; - - double norm_gold; + int type; + int flags; + bool useRoi; virtual void SetUp() { devInfo = GET_PARAM(0); - normCode = GET_PARAM(1); - useRoi = GET_PARAM(2); + size = GET_PARAM(1); + type = GET_PARAM(2); + flags = GET_PARAM(3); + useRoi = GET_PARAM(4); cv::gpu::setDevice(devInfo.deviceID()); - - cv::RNG& rng = TS::ptr()->get_rng(); - - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); - - mat1 = randomMat(rng, size, CV_8UC1, 1, 255, false); - mat2 = randomMat(rng, size, CV_8UC1, 1, 255, false); - - norm_gold = cv::norm(mat1, mat2, normCode); } }; -TEST_P(NormDiff, Accuracy) +TEST_P(GEMM, Accuracy) { - double norm = cv::gpu::norm(loadMat(mat1, useRoi), loadMat(mat2, useRoi), normCode); + cv::Mat src1 = randomMat(size, type, -10.0, 10.0); + cv::Mat src2 = randomMat(size, type, -10.0, 10.0); + cv::Mat src3 = randomMat(size, type, -10.0, 10.0); + double alpha = randomDouble(-10.0, 10.0); + double beta = randomDouble(-10.0, 10.0); + + cv::gpu::GpuMat dst = createMat(size, type, useRoi); + cv::gpu::gemm(loadMat(src1, useRoi), loadMat(src2, useRoi), alpha, loadMat(src3, useRoi), beta, dst, flags); + + cv::Mat dst_gold; + cv::gemm(src1, src2, alpha, src3, beta, dst_gold, flags); - EXPECT_NEAR(norm_gold, norm, 1e-6); + EXPECT_MAT_NEAR(dst_gold, dst, CV_MAT_DEPTH(type) == CV_32F ? 1e-1 : 1e-10); } -INSTANTIATE_TEST_CASE_P(Arithm, NormDiff, Combine( - ALL_DEVICES, - Values((int) cv::NORM_INF, (int) cv::NORM_L1, (int) cv::NORM_L2), - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, GEMM, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_32FC1), MatType(CV_32FC2), MatType(CV_64FC1), MatType(CV_64FC2)), + ALL_GEMM_FLAGS, + WHOLE_SUBMAT)); //////////////////////////////////////////////////////////////////////////////// -// flip +// Transpose -PARAM_TEST_CASE(Flip, cv::gpu::DeviceInfo, MatType, FlipCode, UseRoi) +PARAM_TEST_CASE(Transpose, cv::gpu::DeviceInfo, cv::Size, MatType, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; int type; - int flip_code; bool useRoi; - cv::Size size; - cv::Mat mat; - - cv::Mat dst_gold; - virtual void SetUp() { devInfo = GET_PARAM(0); - type = GET_PARAM(1); - flip_code = GET_PARAM(2); + size = GET_PARAM(1); + type = GET_PARAM(2); useRoi = GET_PARAM(3); cv::gpu::setDevice(devInfo.deviceID()); - - cv::RNG& rng = TS::ptr()->get_rng(); - - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); - - mat = randomMat(rng, size, type, 1, 255, false); - - cv::flip(mat, dst_gold, flip_code); } }; -TEST_P(Flip, Accuracy) +TEST_P(Transpose, Accuracy) { - cv::Mat dst; - - cv::gpu::GpuMat gpu_res; + cv::Mat src = randomMat(size, type); - cv::gpu::flip(loadMat(mat, useRoi), gpu_res, flip_code); + cv::gpu::GpuMat dst = createMat(cv::Size(size.height, size.width), type, useRoi); + cv::gpu::transpose(loadMat(src, useRoi), dst); - gpu_res.download(dst); + cv::Mat dst_gold; + cv::transpose(src, dst_gold); EXPECT_MAT_NEAR(dst_gold, dst, 0.0); } -INSTANTIATE_TEST_CASE_P(Arithm, Flip, Combine( - ALL_DEVICES, - Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_16UC1, CV_16UC3, CV_16UC4, CV_32SC1, CV_32SC3, CV_32SC4, CV_32FC1, CV_32FC3, CV_32FC4), - Values((int)FLIP_BOTH, (int)FLIP_X, (int)FLIP_Y), - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, Transpose, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), + MatType(CV_8UC4), + MatType(CV_16UC2), + MatType(CV_16SC2), + MatType(CV_32SC1), + MatType(CV_32SC2), + MatType(CV_64FC1)), + WHOLE_SUBMAT)); //////////////////////////////////////////////////////////////////////////////// -// LUT +// Flip -PARAM_TEST_CASE(LUT, cv::gpu::DeviceInfo, MatType, UseRoi) +PARAM_TEST_CASE(Flip, cv::gpu::DeviceInfo, cv::Size, MatType, FlipCode, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; int type; + int flip_code; bool useRoi; - cv::Size size; - cv::Mat mat; - cv::Mat lut; - - cv::Mat dst_gold; - virtual void SetUp() { devInfo = GET_PARAM(0); - type = GET_PARAM(1); - useRoi = GET_PARAM(2); + size = GET_PARAM(1); + type = GET_PARAM(2); + flip_code = GET_PARAM(3); + useRoi = GET_PARAM(4); cv::gpu::setDevice(devInfo.deviceID()); - - cv::RNG& rng = TS::ptr()->get_rng(); - - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); - - mat = randomMat(rng, size, type, 1, 255, false); - lut = randomMat(rng, cv::Size(256, 1), CV_8UC1, 100, 200, false); - - cv::LUT(mat, lut, dst_gold); } }; -TEST_P(LUT, Accuracy) +TEST_P(Flip, Accuracy) { - cv::Mat dst; - - cv::gpu::GpuMat gpu_res; + cv::Mat src = randomMat(size, type); - cv::gpu::LUT(loadMat(mat, useRoi), lut, gpu_res); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); + cv::gpu::flip(loadMat(src, useRoi), dst, flip_code); - gpu_res.download(dst); + cv::Mat dst_gold; + cv::flip(src, dst_gold, flip_code); EXPECT_MAT_NEAR(dst_gold, dst, 0.0); } -INSTANTIATE_TEST_CASE_P(Arithm, LUT, Combine( - ALL_DEVICES, - Values(CV_8UC1, CV_8UC3), - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, Flip, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), + MatType(CV_8UC3), + MatType(CV_8UC4), + MatType(CV_16UC1), + MatType(CV_16UC3), + MatType(CV_16UC4), + MatType(CV_32SC1), + MatType(CV_32SC3), + MatType(CV_32SC4), + MatType(CV_32FC1), + MatType(CV_32FC3), + MatType(CV_32FC4)), + ALL_FLIP_CODES, + WHOLE_SUBMAT)); //////////////////////////////////////////////////////////////////////////////// -// pow +// LUT -PARAM_TEST_CASE(Pow, cv::gpu::DeviceInfo, MatType, UseRoi) +PARAM_TEST_CASE(LUT, cv::gpu::DeviceInfo, cv::Size, MatType, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; int type; bool useRoi; - double power; - cv::Size size; - cv::Mat mat; - - cv::Mat dst_gold; - virtual void SetUp() { devInfo = GET_PARAM(0); - type = GET_PARAM(1); - useRoi = GET_PARAM(2); + size = GET_PARAM(1); + type = GET_PARAM(2); + useRoi = GET_PARAM(3); cv::gpu::setDevice(devInfo.deviceID()); + } +}; - cv::RNG& rng = TS::ptr()->get_rng(); - - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); +TEST_P(LUT, OneChannel) +{ + cv::Mat src = randomMat(size, type); + cv::Mat lut = randomMat(cv::Size(256, 1), CV_8UC1); - mat = randomMat(rng, size, type, 0.0, 100.0, false); + cv::gpu::GpuMat dst = createMat(size, CV_MAKE_TYPE(lut.depth(), src.channels())); + cv::gpu::LUT(loadMat(src, useRoi), lut, dst); - if (mat.depth() == CV_32F) - power = rng.uniform(1.2f, 3.f); - else - { - int ipower = rng.uniform(2, 8); - power = (float)ipower; - } + cv::Mat dst_gold; + cv::LUT(src, lut, dst_gold); - cv::pow(mat, power, dst_gold); - } -}; + EXPECT_MAT_NEAR(dst_gold, dst, 0.0); +} -TEST_P(Pow, Accuracy) +TEST_P(LUT, MultiChannel) { - cv::Mat dst; - - cv::gpu::GpuMat gpu_res; + cv::Mat src = randomMat(size, type); + cv::Mat lut = randomMat(cv::Size(256, 1), CV_MAKE_TYPE(CV_8U, src.channels())); - cv::gpu::pow(loadMat(mat, useRoi), power, gpu_res); + cv::gpu::GpuMat dst = createMat(size, CV_MAKE_TYPE(lut.depth(), src.channels()), useRoi); + cv::gpu::LUT(loadMat(src, useRoi), lut, dst); - gpu_res.download(dst); + cv::Mat dst_gold; + cv::LUT(src, lut, dst_gold); - EXPECT_MAT_NEAR(dst_gold, dst, 2); + EXPECT_MAT_NEAR(dst_gold, dst, 0.0); } -INSTANTIATE_TEST_CASE_P(Arithm, Pow, Combine( - ALL_DEVICES, - Values(CV_32F, CV_32FC3), - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, LUT, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), MatType(CV_8UC3)), + WHOLE_SUBMAT)); //////////////////////////////////////////////////////////////////////////////// -// magnitude +// Magnitude -PARAM_TEST_CASE(Magnitude, cv::gpu::DeviceInfo, UseRoi) +PARAM_TEST_CASE(Magnitude, cv::gpu::DeviceInfo, cv::Size, UseRoi) { cv::gpu::DeviceInfo devInfo; - bool useRoi; - cv::Size size; - cv::Mat mat1, mat2; - - cv::Mat dst_gold; + bool useRoi; virtual void SetUp() { devInfo = GET_PARAM(0); - useRoi = GET_PARAM(1); + size = GET_PARAM(1); + useRoi = GET_PARAM(2); cv::gpu::setDevice(devInfo.deviceID()); - - cv::RNG& rng = TS::ptr()->get_rng(); - - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); - - mat1 = randomMat(rng, size, CV_32FC1, 0.0, 100.0, false); - mat2 = randomMat(rng, size, CV_32FC1, 0.0, 100.0, false); - - cv::magnitude(mat1, mat2, dst_gold); } }; -TEST_P(Magnitude, Accuracy) +TEST_P(Magnitude, NPP) { - cv::Mat dst; + cv::Mat src = randomMat(size, CV_32FC2); - cv::gpu::GpuMat gpu_res; + cv::gpu::GpuMat dst = createMat(size, CV_32FC1, useRoi); + cv::gpu::magnitude(loadMat(src, useRoi), dst); - cv::gpu::magnitude(loadMat(mat1, useRoi), loadMat(mat2, useRoi), gpu_res); - - gpu_res.download(dst); + cv::Mat arr[2]; + cv::split(src, arr); + cv::Mat dst_gold; + cv::magnitude(arr[0], arr[1], dst_gold); EXPECT_MAT_NEAR(dst_gold, dst, 1e-4); } -INSTANTIATE_TEST_CASE_P(Arithm, Magnitude, Combine( - ALL_DEVICES, - WHOLE_SUBMAT)); - -//////////////////////////////////////////////////////////////////////////////// -// phase - -PARAM_TEST_CASE(Phase, cv::gpu::DeviceInfo, UseRoi) +TEST_P(Magnitude, Sqr_NPP) { - cv::gpu::DeviceInfo devInfo; - bool useRoi; + cv::Mat src = randomMat(size, CV_32FC2); - cv::Size size; - cv::Mat mat1, mat2; + cv::gpu::GpuMat dst = createMat(size, CV_32FC1, useRoi); + cv::gpu::magnitudeSqr(loadMat(src, useRoi), dst); + cv::Mat arr[2]; + cv::split(src, arr); cv::Mat dst_gold; + cv::magnitude(arr[0], arr[1], dst_gold); + cv::multiply(dst_gold, dst_gold, dst_gold); - virtual void SetUp() - { - devInfo = GET_PARAM(0); - useRoi = GET_PARAM(1); + EXPECT_MAT_NEAR(dst_gold, dst, 1e-1); +} - cv::gpu::setDevice(devInfo.deviceID()); +TEST_P(Magnitude, Accuracy) +{ + cv::Mat x = randomMat(size, CV_32FC1); + cv::Mat y = randomMat(size, CV_32FC1); - cv::RNG& rng = TS::ptr()->get_rng(); + cv::gpu::GpuMat dst = createMat(size, CV_32FC1, useRoi); + cv::gpu::magnitude(loadMat(x, useRoi), loadMat(y, useRoi), dst); - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); - - mat1 = randomMat(rng, size, CV_32FC1, 0.0, 100.0, false); - mat2 = randomMat(rng, size, CV_32FC1, 0.0, 100.0, false); + cv::Mat dst_gold; + cv::magnitude(x, y, dst_gold); - cv::phase(mat1, mat2, dst_gold); - } -}; + EXPECT_MAT_NEAR(dst_gold, dst, 1e-4); +} -TEST_P(Phase, Accuracy) +TEST_P(Magnitude, Sqr_Accuracy) { - cv::Mat dst; + cv::Mat x = randomMat(size, CV_32FC1); + cv::Mat y = randomMat(size, CV_32FC1); - cv::gpu::GpuMat gpu_res; + cv::gpu::GpuMat dst = createMat(size, CV_32FC1, useRoi); + cv::gpu::magnitudeSqr(loadMat(x, useRoi), loadMat(y, useRoi), dst); - cv::gpu::phase(loadMat(mat1, useRoi), loadMat(mat2, useRoi), gpu_res); - - gpu_res.download(dst); + cv::Mat dst_gold; + cv::magnitude(x, y, dst_gold); + cv::multiply(dst_gold, dst_gold, dst_gold); - EXPECT_MAT_NEAR(dst_gold, dst, 1e-3); + EXPECT_MAT_NEAR(dst_gold, dst, 1e-1); } -INSTANTIATE_TEST_CASE_P(Arithm, Phase, Combine( - ALL_DEVICES, - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, Magnitude, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + WHOLE_SUBMAT)); //////////////////////////////////////////////////////////////////////////////// -// cartToPolar +// Phase -PARAM_TEST_CASE(CartToPolar, cv::gpu::DeviceInfo, UseRoi) +PARAM_TEST_CASE(Phase, cv::gpu::DeviceInfo, cv::Size, bool, UseRoi) { cv::gpu::DeviceInfo devInfo; - bool useRoi; - cv::Size size; - cv::Mat mat1, mat2; - - cv::Mat mag_gold; - cv::Mat angle_gold; + bool angleInDegrees; + bool useRoi; virtual void SetUp() { devInfo = GET_PARAM(0); - useRoi = GET_PARAM(1); + size = GET_PARAM(1); + angleInDegrees = GET_PARAM(2); + useRoi = GET_PARAM(3); cv::gpu::setDevice(devInfo.deviceID()); + } +}; + +TEST_P(Phase, Accuracy) +{ + cv::Mat x = randomMat(size, CV_32FC1); + cv::Mat y = randomMat(size, CV_32FC1); + + cv::gpu::GpuMat dst = createMat(size, CV_32FC1, useRoi); + cv::gpu::phase(loadMat(x, useRoi), loadMat(y, useRoi), dst, angleInDegrees); - cv::RNG& rng = TS::ptr()->get_rng(); + cv::Mat dst_gold; + cv::phase(x, y, dst_gold, angleInDegrees); - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); + EXPECT_MAT_NEAR(dst_gold, dst, angleInDegrees ? 1e-2 : 1e-3); +} - mat1 = randomMat(rng, size, CV_32FC1, -100.0, 100.0, false); - mat2 = randomMat(rng, size, CV_32FC1, -100.0, 100.0, false); +INSTANTIATE_TEST_CASE_P(GPU_Core, Phase, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Bool(), + WHOLE_SUBMAT)); - cv::cartToPolar(mat1, mat2, mag_gold, angle_gold); +//////////////////////////////////////////////////////////////////////////////// +// CartToPolar + +PARAM_TEST_CASE(CartToPolar, cv::gpu::DeviceInfo, cv::Size, bool, UseRoi) +{ + cv::gpu::DeviceInfo devInfo; + cv::Size size; + bool angleInDegrees; + bool useRoi; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + angleInDegrees = GET_PARAM(2); + useRoi = GET_PARAM(3); + + cv::gpu::setDevice(devInfo.deviceID()); } }; TEST_P(CartToPolar, Accuracy) { - cv::Mat mag, angle; - - cv::gpu::GpuMat gpuMag; - cv::gpu::GpuMat gpuAngle; + cv::Mat x = randomMat(size, CV_32FC1); + cv::Mat y = randomMat(size, CV_32FC1); - cv::gpu::cartToPolar(loadMat(mat1, useRoi), loadMat(mat2, useRoi), gpuMag, gpuAngle); + cv::gpu::GpuMat mag = createMat(size, CV_32FC1, useRoi); + cv::gpu::GpuMat angle = createMat(size, CV_32FC1, useRoi); + cv::gpu::cartToPolar(loadMat(x, useRoi), loadMat(y, useRoi), mag, angle, angleInDegrees); - gpuMag.download(mag); - gpuAngle.download(angle); + cv::Mat mag_gold; + cv::Mat angle_gold; + cv::cartToPolar(x, y, mag_gold, angle_gold, angleInDegrees); EXPECT_MAT_NEAR(mag_gold, mag, 1e-4); - EXPECT_MAT_NEAR(angle_gold, angle, 1e-3); + EXPECT_MAT_NEAR(angle_gold, angle, angleInDegrees ? 1e-2 : 1e-3); } -INSTANTIATE_TEST_CASE_P(Arithm, CartToPolar, Combine( - ALL_DEVICES, - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, CartToPolar, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Bool(), + WHOLE_SUBMAT)); //////////////////////////////////////////////////////////////////////////////// // polarToCart -PARAM_TEST_CASE(PolarToCart, cv::gpu::DeviceInfo, UseRoi) +PARAM_TEST_CASE(PolarToCart, cv::gpu::DeviceInfo, cv::Size, bool, UseRoi) { cv::gpu::DeviceInfo devInfo; - bool useRoi; - cv::Size size; - cv::Mat mag; - cv::Mat angle; - - cv::Mat x_gold; - cv::Mat y_gold; + bool angleInDegrees; + bool useRoi; virtual void SetUp() { devInfo = GET_PARAM(0); - useRoi = GET_PARAM(1); + size = GET_PARAM(1); + angleInDegrees = GET_PARAM(2); + useRoi = GET_PARAM(3); cv::gpu::setDevice(devInfo.deviceID()); - - cv::RNG& rng = TS::ptr()->get_rng(); - - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); - - mag = randomMat(rng, size, CV_32FC1, -100.0, 100.0, false); - angle = randomMat(rng, size, CV_32FC1, 0.0, 2.0 * CV_PI, false); - - cv::polarToCart(mag, angle, x_gold, y_gold); } }; TEST_P(PolarToCart, Accuracy) { - cv::Mat x, y; - - cv::gpu::GpuMat gpuX; - cv::gpu::GpuMat gpuY; + cv::Mat magnitude = randomMat(size, CV_32FC1); + cv::Mat angle = randomMat(size, CV_32FC1); - cv::gpu::polarToCart(loadMat(mag, useRoi), loadMat(angle, useRoi), gpuX, gpuY); + cv::gpu::GpuMat x = createMat(size, CV_32FC1, useRoi); + cv::gpu::GpuMat y = createMat(size, CV_32FC1, useRoi); + cv::gpu::polarToCart(loadMat(magnitude, useRoi), loadMat(angle, useRoi), x, y, angleInDegrees); - gpuX.download(x); - gpuY.download(y); + cv::Mat x_gold; + cv::Mat y_gold; + cv::polarToCart(magnitude, angle, x_gold, y_gold, angleInDegrees); EXPECT_MAT_NEAR(x_gold, x, 1e-4); EXPECT_MAT_NEAR(y_gold, y, 1e-4); } -INSTANTIATE_TEST_CASE_P(Arithm, PolarToCart, Combine( - ALL_DEVICES, - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, PolarToCart, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Bool(), + WHOLE_SUBMAT)); //////////////////////////////////////////////////////////////////////////////// -// minMax +// MeanStdDev -PARAM_TEST_CASE(MinMax, cv::gpu::DeviceInfo, MatType, UseRoi) +PARAM_TEST_CASE(MeanStdDev, cv::gpu::DeviceInfo, cv::Size, UseRoi) { cv::gpu::DeviceInfo devInfo; - int type; - bool useRoi; - cv::Size size; - cv::Mat mat; - cv::Mat mask; - - double minVal_gold; - double maxVal_gold; + bool useRoi; virtual void SetUp() { devInfo = GET_PARAM(0); - type = GET_PARAM(1); + size = GET_PARAM(1); useRoi = GET_PARAM(2); cv::gpu::setDevice(devInfo.deviceID()); + } +}; - cv::RNG& rng = TS::ptr()->get_rng(); +TEST_P(MeanStdDev, Accuracy) +{ + cv::Mat src = randomMat(size, CV_8UC1); - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); + cv::Scalar mean; + cv::Scalar stddev; + cv::gpu::meanStdDev(loadMat(src, useRoi), mean, stddev); - mat = randomMat(rng, size, type, 0.0, 127.0, false); - mask = randomMat(rng, size, CV_8UC1, 0, 2, false); + cv::Scalar mean_gold; + cv::Scalar stddev_gold; + cv::meanStdDev(src, mean_gold, stddev_gold); - if (type != CV_8S) - { - cv::minMaxLoc(mat, &minVal_gold, &maxVal_gold, 0, 0, mask); - } - else - { - // OpenCV's minMaxLoc doesn't support CV_8S type - minVal_gold = std::numeric_limits::max(); - maxVal_gold = -std::numeric_limits::max(); - for (int i = 0; i < mat.rows; ++i) - { - const signed char* mat_row = mat.ptr(i); - const unsigned char* mask_row = mask.ptr(i); - for (int j = 0; j < mat.cols; ++j) - { - if (mask_row[j]) - { - signed char val = mat_row[j]; - if (val < minVal_gold) minVal_gold = val; - if (val > maxVal_gold) maxVal_gold = val; - } - } - } - } + EXPECT_SCALAR_NEAR(mean_gold, mean, 1e-5); + EXPECT_SCALAR_NEAR(stddev_gold, stddev, 1e-5); +} + +INSTANTIATE_TEST_CASE_P(GPU_Core, MeanStdDev, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + WHOLE_SUBMAT)); + +//////////////////////////////////////////////////////////////////////////////// +// Norm + +PARAM_TEST_CASE(Norm, cv::gpu::DeviceInfo, cv::Size, MatDepth, NormCode, UseRoi) +{ + cv::gpu::DeviceInfo devInfo; + cv::Size size; + int depth; + int normCode; + bool useRoi; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + depth = GET_PARAM(2); + normCode = GET_PARAM(3); + useRoi = GET_PARAM(4); + + cv::gpu::setDevice(devInfo.deviceID()); } }; -TEST_P(MinMax, Accuracy) +TEST_P(Norm, Accuracy) { - if (type == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE)) - return; + cv::Mat src = randomMat(size, depth); - double minVal, maxVal; + double val = cv::gpu::norm(loadMat(src, useRoi), normCode); - cv::gpu::minMax(loadMat(mat, useRoi), &minVal, &maxVal, loadMat(mask, useRoi)); + double val_gold = cv::norm(src, normCode); - EXPECT_DOUBLE_EQ(minVal_gold, minVal); - EXPECT_DOUBLE_EQ(maxVal_gold, maxVal); + EXPECT_NEAR(val_gold, val, depth < CV_32F ? 0.0 : 1.0); } -INSTANTIATE_TEST_CASE_P(Arithm, MinMax, Combine( - ALL_DEVICES, - Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, Norm, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatDepth(CV_8U), + MatDepth(CV_8S), + MatDepth(CV_16U), + MatDepth(CV_16S), + MatDepth(CV_32S), + MatDepth(CV_32F)), + testing::Values(NormCode(cv::NORM_L1), NormCode(cv::NORM_L2), NormCode(cv::NORM_INF)), + WHOLE_SUBMAT)); //////////////////////////////////////////////////////////////////////////////// -// minMaxLoc +// normDiff -PARAM_TEST_CASE(MinMaxLoc, cv::gpu::DeviceInfo, MatType, UseRoi) +PARAM_TEST_CASE(NormDiff, cv::gpu::DeviceInfo, cv::Size, NormCode, UseRoi) { cv::gpu::DeviceInfo devInfo; - int type; - bool useRoi; - cv::Size size; - cv::Mat mat; - cv::Mat mask; - - double minVal_gold; - double maxVal_gold; - cv::Point minLoc_gold; - cv::Point maxLoc_gold; + int normCode; + bool useRoi; virtual void SetUp() { devInfo = GET_PARAM(0); - type = GET_PARAM(1); - useRoi = GET_PARAM(2); + size = GET_PARAM(1); + normCode = GET_PARAM(2); + useRoi = GET_PARAM(3); cv::gpu::setDevice(devInfo.deviceID()); + } +}; - cv::RNG& rng = TS::ptr()->get_rng(); +TEST_P(NormDiff, Accuracy) +{ + cv::Mat src1 = randomMat(size, CV_8UC1); + cv::Mat src2 = randomMat(size, CV_8UC1); - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); + double val = cv::gpu::norm(loadMat(src1, useRoi), loadMat(src2, useRoi), normCode); - mat = randomMat(rng, size, type, 0.0, 127.0, false); - mask = randomMat(rng, size, CV_8UC1, 0, 2, false); + double val_gold = cv::norm(src1, src2, normCode); - if (type != CV_8S) - { - cv::minMaxLoc(mat, &minVal_gold, &maxVal_gold, &minLoc_gold, &maxLoc_gold, mask); - } - else + EXPECT_NEAR(val_gold, val, 0.0); +} + +INSTANTIATE_TEST_CASE_P(GPU_Core, NormDiff, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(NormCode(cv::NORM_L1), NormCode(cv::NORM_L2), NormCode(cv::NORM_INF)), + WHOLE_SUBMAT)); + +////////////////////////////////////////////////////////////////////////////// +// Sum + +namespace +{ + template + cv::Scalar absSumImpl(const cv::Mat& src) + { + const int cn = src.channels(); + + cv::Scalar sum = cv::Scalar::all(0); + + for (int y = 0; y < src.rows; ++y) { - // OpenCV's minMaxLoc doesn't support CV_8S type - minVal_gold = std::numeric_limits::max(); - maxVal_gold = -std::numeric_limits::max(); - for (int i = 0; i < mat.rows; ++i) + for (int x = 0; x < src.cols; ++x) { - const signed char* mat_row = mat.ptr(i); - const unsigned char* mask_row = mask.ptr(i); - for (int j = 0; j < mat.cols; ++j) - { - if (mask_row[j]) - { - signed char val = mat_row[j]; - if (val < minVal_gold) { minVal_gold = val; minLoc_gold = cv::Point(j, i); } - if (val > maxVal_gold) { maxVal_gold = val; maxLoc_gold = cv::Point(j, i); } - } - } + for (int c = 0; c < cn; ++c) + sum[c] += std::abs(src.at(y, x * cn + c)); } } + + return sum; } -}; -TEST_P(MinMaxLoc, Accuracy) -{ - if (type == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE)) - return; + cv::Scalar absSumGold(const cv::Mat& src) + { + typedef cv::Scalar (*func_t)(const cv::Mat& src); - double minVal, maxVal; - cv::Point minLoc, maxLoc; + static const func_t funcs[] = + { + absSumImpl, + absSumImpl, + absSumImpl, + absSumImpl, + absSumImpl, + absSumImpl, + absSumImpl + }; - cv::gpu::minMaxLoc(loadMat(mat, useRoi), &minVal, &maxVal, &minLoc, &maxLoc, loadMat(mask, useRoi)); + return funcs[src.depth()](src); + } - EXPECT_DOUBLE_EQ(minVal_gold, minVal); - EXPECT_DOUBLE_EQ(maxVal_gold, maxVal); + template + cv::Scalar sqrSumImpl(const cv::Mat& src) + { + const int cn = src.channels(); - int cmpMinVals = memcmp(mat.data + minLoc_gold.y * mat.step + minLoc_gold.x * mat.elemSize(), - mat.data + minLoc.y * mat.step + minLoc.x * mat.elemSize(), - mat.elemSize()); - int cmpMaxVals = memcmp(mat.data + maxLoc_gold.y * mat.step + maxLoc_gold.x * mat.elemSize(), - mat.data + maxLoc.y * mat.step + maxLoc.x * mat.elemSize(), - mat.elemSize()); + cv::Scalar sum = cv::Scalar::all(0); - EXPECT_EQ(0, cmpMinVals); - EXPECT_EQ(0, cmpMaxVals); -} + for (int y = 0; y < src.rows; ++y) + { + for (int x = 0; x < src.cols; ++x) + { + for (int c = 0; c < cn; ++c) + { + const T val = src.at(y, x * cn + c); + sum[c] += val * val; + } + } + } + + return sum; + } -INSTANTIATE_TEST_CASE_P(Arithm, MinMaxLoc, Combine( - ALL_DEVICES, - Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), - WHOLE_SUBMAT)); + cv::Scalar sqrSumGold(const cv::Mat& src) + { + typedef cv::Scalar (*func_t)(const cv::Mat& src); -//////////////////////////////////////////////////////////////////////////// -// countNonZero + static const func_t funcs[] = + { + sqrSumImpl, + sqrSumImpl, + sqrSumImpl, + sqrSumImpl, + sqrSumImpl, + sqrSumImpl, + sqrSumImpl + }; + + return funcs[src.depth()](src); + } +} -PARAM_TEST_CASE(CountNonZero, cv::gpu::DeviceInfo, MatType, UseRoi) +PARAM_TEST_CASE(Sum, cv::gpu::DeviceInfo, cv::Size, MatType, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; int type; bool useRoi; - cv::Size size; - cv::Mat mat; - - int n_gold; + cv::Mat src; virtual void SetUp() { devInfo = GET_PARAM(0); - type = GET_PARAM(1); - useRoi = GET_PARAM(2); + size = GET_PARAM(1); + type = GET_PARAM(2); + useRoi = GET_PARAM(3); cv::gpu::setDevice(devInfo.deviceID()); - cv::RNG& rng = TS::ptr()->get_rng(); + src = randomMat(size, type, -128.0, 128.0); + } +}; - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); +TEST_P(Sum, Simple) +{ + cv::Scalar val = cv::gpu::sum(loadMat(src, useRoi)); - cv::Mat matBase = randomMat(rng, size, CV_8U, 0.0, 1.0, false); - matBase.convertTo(mat, type); + cv::Scalar val_gold = cv::sum(src); - n_gold = cv::countNonZero(mat); - } -}; + EXPECT_SCALAR_NEAR(val_gold, val, CV_MAT_DEPTH(type) < CV_32F ? 0.0 : 0.5); +} -TEST_P(CountNonZero, Accuracy) +TEST_P(Sum, Abs) { - if (type == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE)) - return; + cv::Scalar val = cv::gpu::absSum(loadMat(src, useRoi)); - int n = cv::gpu::countNonZero(loadMat(mat, useRoi)); + cv::Scalar val_gold = absSumGold(src); - ASSERT_EQ(n_gold, n); + EXPECT_SCALAR_NEAR(val_gold, val, CV_MAT_DEPTH(type) < CV_32F ? 0.0 : 0.5); } -INSTANTIATE_TEST_CASE_P(Arithm, CountNonZero, Combine( - ALL_DEVICES, - Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), - WHOLE_SUBMAT)); +TEST_P(Sum, Sqr) +{ + cv::Scalar val = cv::gpu::sqrSum(loadMat(src, useRoi)); -////////////////////////////////////////////////////////////////////////////// -// sum + cv::Scalar val_gold = sqrSumGold(src); + + EXPECT_SCALAR_NEAR(val_gold, val, CV_MAT_DEPTH(type) < CV_32F ? 0.0 : 10); +} + +INSTANTIATE_TEST_CASE_P(GPU_Core, Sum, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + TYPES(CV_8U, CV_32F, 1, 4), + WHOLE_SUBMAT)); -PARAM_TEST_CASE(Sum, cv::gpu::DeviceInfo, MatType, UseRoi) +//////////////////////////////////////////////////////////////////////////////// +// MinMax + +namespace { - cv::gpu::DeviceInfo devInfo; - int type; - bool useRoi; + void minMaxLocGold(const cv::Mat& src, double* minVal_, double* maxVal_ = 0, cv::Point* minLoc_ = 0, cv::Point* maxLoc_ = 0, const cv::Mat& mask = cv::Mat()) + { + if (src.depth() != CV_8S) + { + cv::minMaxLoc(src, minVal_, maxVal_, minLoc_, maxLoc_, mask); + return; + } + + // OpenCV's minMaxLoc doesn't support CV_8S type + double minVal = std::numeric_limits::max(); + cv::Point minLoc(-1, -1); + + double maxVal = -std::numeric_limits::max(); + cv::Point maxLoc(-1, -1); + + for (int y = 0; y < src.rows; ++y) + { + const schar* src_row = src.ptr(y); + const uchar* mask_row = mask.empty() ? 0 : mask.ptr(y); + + for (int x = 0; x < src.cols; ++x) + { + if (!mask_row || mask_row[x]) + { + schar val = src_row[x]; + + if (val < minVal) + { + minVal = val; + minLoc = cv::Point(x, y); + } + + if (val > maxVal) + { + maxVal = val; + maxLoc = cv::Point(x, y); + } + } + } + } + if (minVal_) *minVal_ = minVal; + if (maxVal_) *maxVal_ = maxVal; + + if (minLoc_) *minLoc_ = minLoc; + if (maxLoc_) *maxLoc_ = maxLoc; + } +} + +PARAM_TEST_CASE(MinMax, cv::gpu::DeviceInfo, cv::Size, MatDepth, UseRoi) +{ + cv::gpu::DeviceInfo devInfo; cv::Size size; - cv::Mat mat; + int depth; + bool useRoi; virtual void SetUp() { devInfo = GET_PARAM(0); - type = GET_PARAM(1); - useRoi = GET_PARAM(2); + size = GET_PARAM(1); + depth = GET_PARAM(2); + useRoi = GET_PARAM(3); cv::gpu::setDevice(devInfo.deviceID()); - - cv::RNG& rng = TS::ptr()->get_rng(); - - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); - - mat = randomMat(rng, size, CV_8U, 0.0, 10.0, false); } }; -TEST_P(Sum, Simple) +TEST_P(MinMax, WithoutMask) { - if (type == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE)) - return; + cv::Mat src = randomMat(size, depth); - cv::Scalar sum_gold = cv::sum(mat); + double minVal, maxVal; + cv::gpu::minMax(loadMat(src, useRoi), &minVal, &maxVal); - cv::Scalar sum = cv::gpu::sum(loadMat(mat, useRoi)); + double minVal_gold, maxVal_gold; + minMaxLocGold(src, &minVal_gold, &maxVal_gold); - EXPECT_NEAR(sum[0], sum_gold[0], mat.size().area() * 1e-5); - EXPECT_NEAR(sum[1], sum_gold[1], mat.size().area() * 1e-5); - EXPECT_NEAR(sum[2], sum_gold[2], mat.size().area() * 1e-5); - EXPECT_NEAR(sum[3], sum_gold[3], mat.size().area() * 1e-5); + EXPECT_DOUBLE_EQ(minVal_gold, minVal); + EXPECT_DOUBLE_EQ(maxVal_gold, maxVal); } -TEST_P(Sum, Abs) +TEST_P(MinMax, WithMask) { - if (type == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE)) - return; + cv::Mat src = randomMat(size, depth); + cv::Mat mask = randomMat(size, CV_8UC1, 0.0, 2.0); - cv::Scalar sum_gold = cv::norm(mat, cv::NORM_L1); + double minVal, maxVal; + cv::gpu::minMax(loadMat(src, useRoi), &minVal, &maxVal, loadMat(mask, useRoi)); - cv::Scalar sum = cv::gpu::absSum(loadMat(mat, useRoi)); + double minVal_gold, maxVal_gold; + minMaxLocGold(src, &minVal_gold, &maxVal_gold, 0, 0, mask); - EXPECT_NEAR(sum[0], sum_gold[0], mat.size().area() * 1e-5); - EXPECT_NEAR(sum[1], sum_gold[1], mat.size().area() * 1e-5); - EXPECT_NEAR(sum[2], sum_gold[2], mat.size().area() * 1e-5); - EXPECT_NEAR(sum[3], sum_gold[3], mat.size().area() * 1e-5); + EXPECT_DOUBLE_EQ(minVal_gold, minVal); + EXPECT_DOUBLE_EQ(maxVal_gold, maxVal); } -TEST_P(Sum, Sqr) +TEST_P(MinMax, NullPtr) { - if (type == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE)) - return; + cv::Mat src = randomMat(size, depth); - cv::Mat sqrmat; - multiply(mat, mat, sqrmat); - cv::Scalar sum_gold = sum(sqrmat); + cv::gpu::minMax(loadMat(src, useRoi), 0, 0); +} + +INSTANTIATE_TEST_CASE_P(GPU_Core, MinMax, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + ALL_DEPTH, + WHOLE_SUBMAT)); - cv::Scalar sum = cv::gpu::sqrSum(loadMat(mat, useRoi)); +//////////////////////////////////////////////////////////////////////////////// +// MinMaxLoc - EXPECT_NEAR(sum[0], sum_gold[0], mat.size().area() * 1e-5); - EXPECT_NEAR(sum[1], sum_gold[1], mat.size().area() * 1e-5); - EXPECT_NEAR(sum[2], sum_gold[2], mat.size().area() * 1e-5); - EXPECT_NEAR(sum[3], sum_gold[3], mat.size().area() * 1e-5); -} +namespace +{ + template + void expectEqualImpl(const cv::Mat& src, cv::Point loc_gold, cv::Point loc) + { + EXPECT_EQ(src.at(loc_gold.y, loc_gold.x), src.at(loc.y, loc.x)); + } -INSTANTIATE_TEST_CASE_P(Arithm, Sum, Combine( - ALL_DEVICES, - Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), - WHOLE_SUBMAT)); + void expectEqual(const cv::Mat& src, cv::Point loc_gold, cv::Point loc) + { + typedef void (*func_t)(const cv::Mat& src, cv::Point loc_gold, cv::Point loc); -////////////////////////////////////////////////////////////////////////////// -// addWeighted + static const func_t funcs[] = + { + expectEqualImpl, + expectEqualImpl, + expectEqualImpl, + expectEqualImpl, + expectEqualImpl, + expectEqualImpl, + expectEqualImpl + }; + + funcs[src.depth()](src, loc_gold, loc); + } +} -PARAM_TEST_CASE(AddWeighted, cv::gpu::DeviceInfo, MatType, MatType, MatType, UseRoi) +PARAM_TEST_CASE(MinMaxLoc, cv::gpu::DeviceInfo, cv::Size, MatDepth, UseRoi) { cv::gpu::DeviceInfo devInfo; - int type1; - int type2; - int dtype; - bool useRoi; - cv::Size size; - cv::Mat src1; - cv::Mat src2; - double alpha; - double beta; - double gamma; - - cv::Mat dst_gold; + int depth; + bool useRoi; virtual void SetUp() { devInfo = GET_PARAM(0); - type1 = GET_PARAM(1); - type2 = GET_PARAM(2); - dtype = GET_PARAM(3); - useRoi = GET_PARAM(4); + size = GET_PARAM(1); + depth = GET_PARAM(2); + useRoi = GET_PARAM(3); cv::gpu::setDevice(devInfo.deviceID()); + } +}; - cv::RNG& rng = TS::ptr()->get_rng(); +TEST_P(MinMaxLoc, WithoutMask) +{ + cv::Mat src = randomMat(size, depth); - size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200)); + double minVal, maxVal; + cv::Point minLoc, maxLoc; + cv::gpu::minMaxLoc(loadMat(src, useRoi), &minVal, &maxVal, &minLoc, &maxLoc); - src1 = randomMat(rng, size, type1, 0.0, 255.0, false); - src2 = randomMat(rng, size, type2, 0.0, 255.0, false); + double minVal_gold, maxVal_gold; + cv::Point minLoc_gold, maxLoc_gold; + minMaxLocGold(src, &minVal_gold, &maxVal_gold, &minLoc_gold, &maxLoc_gold); - alpha = rng.uniform(-10.0, 10.0); - beta = rng.uniform(-10.0, 10.0); - gamma = rng.uniform(-10.0, 10.0); + EXPECT_DOUBLE_EQ(minVal_gold, minVal); + EXPECT_DOUBLE_EQ(maxVal_gold, maxVal); - cv::addWeighted(src1, alpha, src2, beta, gamma, dst_gold, dtype); - } -}; + expectEqual(src, minLoc_gold, minLoc); + expectEqual(src, maxLoc_gold, maxLoc); +} -TEST_P(AddWeighted, Accuracy) +TEST_P(MinMaxLoc, WithMask) { - if ((src1.depth() == CV_64F || src2.depth() == CV_64F || dst_gold.depth() == CV_64F) && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE)) - return; + cv::Mat src = randomMat(size, depth); + cv::Mat mask = randomMat(size, CV_8UC1, 0.0, 2.0); - cv::Mat dst; + double minVal, maxVal; + cv::Point minLoc, maxLoc; + cv::gpu::minMaxLoc(loadMat(src, useRoi), &minVal, &maxVal, &minLoc, &maxLoc, loadMat(mask, useRoi)); + + double minVal_gold, maxVal_gold; + cv::Point minLoc_gold, maxLoc_gold; + minMaxLocGold(src, &minVal_gold, &maxVal_gold, &minLoc_gold, &maxLoc_gold, mask); - cv::gpu::GpuMat dev_dst; + EXPECT_DOUBLE_EQ(minVal_gold, minVal); + EXPECT_DOUBLE_EQ(maxVal_gold, maxVal); - cv::gpu::addWeighted(loadMat(src1, useRoi), alpha, loadMat(src2, useRoi), beta, gamma, dev_dst, dtype); + expectEqual(src, minLoc_gold, minLoc); + expectEqual(src, maxLoc_gold, maxLoc); +} - dev_dst.download(dst); +TEST_P(MinMaxLoc, NullPtr) +{ + cv::Mat src = randomMat(size, depth); - EXPECT_MAT_NEAR(dst_gold, dst, dtype < CV_32F ? 1.0 : 1e-12); + cv::gpu::minMaxLoc(loadMat(src, useRoi), 0, 0, 0, 0); } -INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, Combine( - ALL_DEVICES, - TYPES(CV_8U, CV_64F, 1, 1), - TYPES(CV_8U, CV_64F, 1, 1), - TYPES(CV_8U, CV_64F, 1, 1), - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, MinMaxLoc, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + ALL_DEPTH, + WHOLE_SUBMAT)); -////////////////////////////////////////////////////////////////////////////// -// reduce +//////////////////////////////////////////////////////////////////////////// +// CountNonZero -PARAM_TEST_CASE(Reduce, cv::gpu::DeviceInfo, MatType, int, ReduceOp, UseRoi) +PARAM_TEST_CASE(CountNonZero, cv::gpu::DeviceInfo, cv::Size, MatDepth, UseRoi) { cv::gpu::DeviceInfo devInfo; - int type; - int dim; - int reduceOp; - bool useRoi; - cv::Size size; - cv::Mat src; + int depth; + bool useRoi; - cv::Mat dst_gold; virtual void SetUp() { devInfo = GET_PARAM(0); - type = GET_PARAM(1); - dim = GET_PARAM(2); - reduceOp = GET_PARAM(3); - useRoi = GET_PARAM(4); + size = GET_PARAM(1); + depth = GET_PARAM(2); + useRoi = GET_PARAM(3); cv::gpu::setDevice(devInfo.deviceID()); - - cv::RNG& rng = TS::ptr()->get_rng(); - - size = cv::Size(rng.uniform(100, 400), rng.uniform(100, 400)); - - src = randomMat(rng, size, type, 0.0, 255.0, false); - - cv::reduce(src, dst_gold, dim, reduceOp, reduceOp == CV_REDUCE_SUM || reduceOp == CV_REDUCE_AVG ? CV_32F : CV_MAT_DEPTH(type)); - - if (dim == 1) - { - dst_gold.cols = dst_gold.rows; - dst_gold.rows = 1; - dst_gold.step = dst_gold.cols * dst_gold.elemSize(); - } } }; -TEST_P(Reduce, Accuracy) +TEST_P(CountNonZero, Accuracy) { - cv::Mat dst; + cv::Mat srcBase = randomMat(size, CV_8U, 0.0, 1.5); + cv::Mat src; + srcBase.convertTo(src, depth); - cv::gpu::GpuMat dev_dst; + int val = cv::gpu::countNonZero(loadMat(src, useRoi)); - cv::gpu::reduce(loadMat(src, useRoi), dev_dst, dim, reduceOp, reduceOp == CV_REDUCE_SUM || reduceOp == CV_REDUCE_AVG ? CV_32F : CV_MAT_DEPTH(type)); + int val_gold = cv::countNonZero(src); - dev_dst.download(dst); - double norm = reduceOp == CV_REDUCE_SUM || reduceOp == CV_REDUCE_AVG ? 1e-1 : 0.0; - EXPECT_MAT_NEAR(dst_gold, dst, norm); + ASSERT_EQ(val_gold, val); } -INSTANTIATE_TEST_CASE_P(Arithm, Reduce, Combine( - ALL_DEVICES, - Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_16UC1, CV_16UC3, CV_16UC4, CV_32FC1, CV_32FC3, CV_32FC4), - Values(0, 1), - Values((int)CV_REDUCE_SUM, (int)CV_REDUCE_AVG, (int)CV_REDUCE_MAX, (int)CV_REDUCE_MIN), - WHOLE_SUBMAT)); +INSTANTIATE_TEST_CASE_P(GPU_Core, CountNonZero, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + ALL_DEPTH, + WHOLE_SUBMAT)); ////////////////////////////////////////////////////////////////////////////// -// gemm +// Reduce -PARAM_TEST_CASE(GEMM, cv::gpu::DeviceInfo, MatType, GemmFlags, UseRoi) +PARAM_TEST_CASE(Reduce, cv::gpu::DeviceInfo, cv::Size, MatDepth, int, ReduceCode, UseRoi) { cv::gpu::DeviceInfo devInfo; - int type; - int flags; + cv::Size size; + int depth; + int channels; + int reduceOp; bool useRoi; - int size; - cv::Mat src1; - cv::Mat src2; - cv::Mat src3; - double alpha; - double beta; - - cv::Mat dst_gold; + int type; + int dst_depth; + int dst_type; virtual void SetUp() { devInfo = GET_PARAM(0); - type = GET_PARAM(1); - flags = GET_PARAM(2); - useRoi = GET_PARAM(3); + size = GET_PARAM(1); + depth = GET_PARAM(2); + channels = GET_PARAM(3); + reduceOp = GET_PARAM(4); + useRoi = GET_PARAM(5); cv::gpu::setDevice(devInfo.deviceID()); - cv::RNG& rng = TS::ptr()->get_rng(); - - size = rng.uniform(100, 200); - - src1 = randomMat(rng, cv::Size(size, size), type, -10.0, 10.0, false); - src2 = randomMat(rng, cv::Size(size, size), type, -10.0, 10.0, false); - src3 = randomMat(rng, cv::Size(size, size), type, -10.0, 10.0, false); - alpha = rng.uniform(-10.0, 10.0); - beta = rng.uniform(-10.0, 10.0); - - cv::gemm(src1, src2, alpha, src3, beta, dst_gold, flags); + type = CV_MAKE_TYPE(depth, channels); + dst_depth = (reduceOp == CV_REDUCE_MAX || reduceOp == CV_REDUCE_MIN) ? depth : CV_32F; + dst_type = CV_MAKE_TYPE(dst_depth, channels); } }; -TEST_P(GEMM, Accuracy) +TEST_P(Reduce, Rows) { - cv::Mat dst; - - cv::gpu::GpuMat dev_dst; + cv::Mat src = randomMat(size, type); - cv::gpu::gemm(loadMat(src1, useRoi), loadMat(src2, useRoi), alpha, loadMat(src3, useRoi), beta, dev_dst, flags); + cv::gpu::GpuMat dst = createMat(cv::Size(src.cols, 1), dst_type, useRoi); + cv::gpu::reduce(loadMat(src, useRoi), dst, 0, reduceOp, dst_depth); - dev_dst.download(dst); + cv::Mat dst_gold; + cv::reduce(src, dst_gold, 0, reduceOp, dst_depth); - EXPECT_MAT_NEAR(dst_gold, dst, 1e-1); + EXPECT_MAT_NEAR(dst_gold, dst, dst_depth < CV_32F ? 0.0 : 1e-2); } -INSTANTIATE_TEST_CASE_P(Arithm, GEMM, Combine( - ALL_DEVICES, - Values(CV_32FC1, CV_32FC2), - Values(0, (int) cv::GEMM_1_T, (int) cv::GEMM_2_T, (int) cv::GEMM_3_T), - WHOLE_SUBMAT)); +TEST_P(Reduce, Cols) +{ + cv::Mat src = randomMat(size, type); + + cv::gpu::GpuMat dst = createMat(cv::Size(src.rows, 1), dst_type, useRoi); + cv::gpu::reduce(loadMat(src, useRoi), dst, 1, reduceOp, dst_depth); + + cv::Mat dst_gold; + cv::reduce(src, dst_gold, 1, reduceOp, dst_depth); + dst_gold.cols = dst_gold.rows; + dst_gold.rows = 1; + dst_gold.step = dst_gold.cols * dst_gold.elemSize(); + + EXPECT_MAT_NEAR(dst_gold, dst, dst_depth < CV_32F ? 0.0 : 1e-2); +} -#endif // HAVE_CUDA +INSTANTIATE_TEST_CASE_P(GPU_Core, Reduce, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatDepth(CV_8U), + MatDepth(CV_16U), + MatDepth(CV_16S), + MatDepth(CV_32F)), + testing::Values(1, 2, 3, 4), + ALL_REDUCE_CODES, + WHOLE_SUBMAT)); diff --git a/modules/gpu/test/test_imgproc.cpp b/modules/gpu/test/test_imgproc.cpp index 231fa2c59b..01baa185f8 100644 --- a/modules/gpu/test/test_imgproc.cpp +++ b/modules/gpu/test/test_imgproc.cpp @@ -2362,53 +2362,6 @@ TEST_P(ColumnSum, Accuracy) INSTANTIATE_TEST_CASE_P(ImgProc, ColumnSum, ALL_DEVICES); -//////////////////////////////////////////////////////////////////////// -// Norm - -PARAM_TEST_CASE(Norm, cv::gpu::DeviceInfo, MatType, NormCode, UseRoi) -{ - cv::gpu::DeviceInfo devInfo; - int type; - int normType; - bool useRoi; - - cv::Size size; - cv::Mat src; - - double gold; - - virtual void SetUp() - { - devInfo = GET_PARAM(0); - type = GET_PARAM(1); - normType = GET_PARAM(2); - useRoi = GET_PARAM(3); - - cv::gpu::setDevice(devInfo.deviceID()); - - cv::RNG& rng = TS::ptr()->get_rng(); - - size = cv::Size(rng.uniform(100, 400), rng.uniform(100, 400)); - - src = randomMat(rng, size, type, 0.0, 10.0, false); - - gold = cv::norm(src, normType); - } -}; - -TEST_P(Norm, Accuracy) -{ - double res = cv::gpu::norm(loadMat(src, useRoi), normType); - - ASSERT_NEAR(res, gold, 0.5); -} - -INSTANTIATE_TEST_CASE_P(ImgProc, Norm, Combine( - ALL_DEVICES, - TYPES(CV_8U, CV_32F, 1, 1), - Values((int) cv::NORM_INF, (int) cv::NORM_L1, (int) cv::NORM_L2), - WHOLE_SUBMAT)); - //////////////////////////////////////////////////////////////////////////////// // reprojectImageTo3D diff --git a/modules/gpu/test/test_threshold.cpp b/modules/gpu/test/test_threshold.cpp index 70473b7c66..23e29ff575 100644 --- a/modules/gpu/test/test_threshold.cpp +++ b/modules/gpu/test/test_threshold.cpp @@ -82,7 +82,7 @@ INSTANTIATE_TEST_CASE_P(GPU_ImgProc, Threshold, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES, testing::Values(MatType(CV_8UC1), MatType(CV_16SC1), MatType(CV_32FC1)), - testing::Values(ThreshOp(cv::THRESH_BINARY), ThreshOp(cv::THRESH_BINARY_INV), ThreshOp(cv::THRESH_TRUNC), ThreshOp(cv::THRESH_TOZERO), ThreshOp(cv::THRESH_TOZERO_INV)), + ALL_THRESH_OPS, WHOLE_SUBMAT)); #endif // HAVE_CUDA diff --git a/modules/gpu/test/utility.cpp b/modules/gpu/test/utility.cpp index 170f4ba724..6fc7915159 100644 --- a/modules/gpu/test/utility.cpp +++ b/modules/gpu/test/utility.cpp @@ -45,6 +45,7 @@ using namespace std; using namespace cv; using namespace cv::gpu; using namespace cvtest; +using namespace testing; int randomInt(int minVal, int maxVal) { @@ -82,9 +83,9 @@ cv::gpu::GpuMat createMat(cv::Size size, int type, bool useRoi) size0.width += randomInt(5, 15); size0.height += randomInt(5, 15); } - + GpuMat d_m(size0, type); - + if (size0 != size) d_m = d_m(Rect((size0.width - size.width) / 2, (size0.height - size.height) / 2, size.width, size.height)); @@ -98,35 +99,6 @@ GpuMat loadMat(const Mat& m, bool useRoi) return d_m; } -void showDiff(InputArray gold_, InputArray actual_, double eps) -{ - Mat gold; - if (gold_.kind() == _InputArray::MAT) - gold = gold_.getMat(); - else - gold_.getGpuMat().download(gold); - - Mat actual; - if (actual_.kind() == _InputArray::MAT) - actual = actual_.getMat(); - else - actual_.getGpuMat().download(actual); - - Mat diff; - absdiff(gold, actual, diff); - threshold(diff, diff, eps, 255.0, cv::THRESH_BINARY); - - namedWindow("gold", WINDOW_NORMAL); - namedWindow("actual", WINDOW_NORMAL); - namedWindow("diff", WINDOW_NORMAL); - - imshow("gold", gold); - imshow("actual", actual); - imshow("diff", diff); - - waitKey(); -} - bool supportFeature(const DeviceInfo& info, FeatureSet feature) { return TargetArchs::builtWith(feature) && info.supports(feature); @@ -159,7 +131,7 @@ const vector& devices() vector devices(FeatureSet feature) { const vector& d = devices(); - + vector devs_filtered; if (TargetArchs::builtWith(feature)) @@ -220,20 +192,50 @@ Mat readImageType(const string& fname, int type) return src; } -double checkNorm(const Mat& m) +namespace { - return norm(m, NORM_INF); + Mat getMat(InputArray arr) + { + if (arr.kind() == _InputArray::GPU_MAT) + { + Mat m; + arr.getGpuMat().download(m); + return m; + } + + return arr.getMat(); + } +} + +void showDiff(InputArray gold_, InputArray actual_, double eps) +{ + Mat gold = getMat(gold_); + Mat actual = getMat(actual_); + + Mat diff; + absdiff(gold, actual, diff); + threshold(diff, diff, eps, 255.0, cv::THRESH_BINARY); + + namedWindow("gold", WINDOW_NORMAL); + namedWindow("actual", WINDOW_NORMAL); + namedWindow("diff", WINDOW_NORMAL); + + imshow("gold", gold); + imshow("actual", actual); + imshow("diff", diff); + + waitKey(); } -double checkNorm(const Mat& m1, const Mat& m2) +double checkNorm(InputArray m1, const InputArray m2) { - return norm(m1, m2, NORM_INF); + return norm(getMat(m1), getMat(m2), NORM_INF); } -double checkSimilarity(const Mat& m1, const Mat& m2) +double checkSimilarity(InputArray m1, InputArray m2) { Mat diff; - matchTemplate(m1, m2, diff, CV_TM_CCORR_NORMED); + matchTemplate(getMat(m1), getMat(m2), diff, CV_TM_CCORR_NORMED); return std::abs(diff.at(0, 0) - 1.f); } diff --git a/modules/gpu/test/utility.hpp b/modules/gpu/test/utility.hpp index c551ceed5b..bf4849e38f 100644 --- a/modules/gpu/test/utility.hpp +++ b/modules/gpu/test/utility.hpp @@ -65,27 +65,30 @@ std::vector devices(cv::gpu::FeatureSet feature); cv::Mat readImage(const std::string& fileName, int flags = cv::IMREAD_COLOR); cv::Mat readImageType(const std::string& fname, int type); -double checkNorm(const cv::Mat& m); -double checkNorm(const cv::Mat& m1, const cv::Mat& m2); -double checkSimilarity(const cv::Mat& m1, const cv::Mat& m2); - -#define EXPECT_MAT_NORM(mat, eps) \ - { \ - EXPECT_LE(checkNorm(cv::Mat(mat)), eps) \ - } +double checkNorm(cv::InputArray m1, cv::InputArray m2); #define EXPECT_MAT_NEAR(mat1, mat2, eps) \ { \ ASSERT_EQ(mat1.type(), mat2.type()); \ ASSERT_EQ(mat1.size(), mat2.size()); \ - EXPECT_LE(checkNorm(cv::Mat(mat1), cv::Mat(mat2)), eps); \ + EXPECT_LE(checkNorm(mat1, mat2), eps); \ + } + +#define EXPECT_SCALAR_NEAR(s1, s2, eps) \ + { \ + EXPECT_NEAR(s1[0], s2[0], eps); \ + EXPECT_NEAR(s1[1], s2[1], eps); \ + EXPECT_NEAR(s1[2], s2[2], eps); \ + EXPECT_NEAR(s1[3], s2[3], eps); \ } +double checkSimilarity(cv::InputArray m1, cv::InputArray m2); + #define EXPECT_MAT_SIMILAR(mat1, mat2, eps) \ { \ ASSERT_EQ(mat1.type(), mat2.type()); \ ASSERT_EQ(mat1.size(), mat2.size()); \ - EXPECT_LE(checkSimilarity(cv::Mat(mat1), cv::Mat(mat2)), eps); \ + EXPECT_LE(checkSimilarity(mat1, mat2), eps); \ } namespace cv { namespace gpu @@ -112,8 +115,10 @@ public: private: bool val_; }; - void PrintTo(const UseRoi& useRoi, std::ostream* os); +#define WHOLE testing::Values(UseRoi(false)) +#define SUBMAT testing::Values(UseRoi(true)) +#define WHOLE_SUBMAT testing::Values(UseRoi(false), UseRoi(true)) class Inverse { @@ -125,25 +130,30 @@ public: private: bool val_; }; - void PrintTo(const Inverse& useRoi, std::ostream* os); +#define DIRECT_INVERSE testing::Values(Inverse(false), Inverse(true)) CV_ENUM(CmpCode, cv::CMP_EQ, cv::CMP_GT, cv::CMP_GE, cv::CMP_LT, cv::CMP_LE, cv::CMP_NE) +#define ALL_CMP_CODES testing::Values(CmpCode(cv::CMP_EQ), CmpCode(cv::CMP_NE), CmpCode(cv::CMP_GT), CmpCode(cv::CMP_GE), CmpCode(cv::CMP_LT), CmpCode(cv::CMP_LE)) CV_ENUM(NormCode, cv::NORM_INF, cv::NORM_L1, cv::NORM_L2, cv::NORM_TYPE_MASK, cv::NORM_RELATIVE, cv::NORM_MINMAX) enum {FLIP_BOTH = 0, FLIP_X = 1, FLIP_Y = -1}; CV_ENUM(FlipCode, FLIP_BOTH, FLIP_X, FLIP_Y) +#define ALL_FLIP_CODES testing::Values(FlipCode(FLIP_BOTH), FlipCode(FLIP_X), FlipCode(FLIP_Y)) -CV_ENUM(ReduceOp, CV_REDUCE_SUM, CV_REDUCE_AVG, CV_REDUCE_MAX, CV_REDUCE_MIN) +CV_ENUM(ReduceCode, CV_REDUCE_SUM, CV_REDUCE_AVG, CV_REDUCE_MAX, CV_REDUCE_MIN) +#define ALL_REDUCE_CODES testing::Values(ReduceCode(CV_REDUCE_SUM), ReduceCode(CV_REDUCE_AVG), ReduceCode(CV_REDUCE_MAX), ReduceCode(CV_REDUCE_MIN)) -CV_FLAGS(GemmFlags, cv::GEMM_1_T, cv::GEMM_2_T, cv::GEMM_3_T); +CV_FLAGS(GemmFlags, 0, cv::GEMM_1_T, cv::GEMM_2_T, cv::GEMM_3_T); +#define ALL_GEMM_FLAGS testing::Values(GemmFlags(0), GemmFlags(cv::GEMM_1_T), GemmFlags(cv::GEMM_2_T), GemmFlags(cv::GEMM_3_T), GemmFlags(cv::GEMM_1_T | cv::GEMM_2_T), GemmFlags(cv::GEMM_1_T | cv::GEMM_3_T), GemmFlags(cv::GEMM_1_T | cv::GEMM_2_T | cv::GEMM_3_T)) CV_ENUM(DistType, cv::gpu::BruteForceMatcher_GPU_base::L1Dist, cv::gpu::BruteForceMatcher_GPU_base::L2Dist) CV_ENUM(MorphOp, cv::MORPH_OPEN, cv::MORPH_CLOSE, cv::MORPH_GRADIENT, cv::MORPH_TOPHAT, cv::MORPH_BLACKHAT) CV_ENUM(ThreshOp, cv::THRESH_BINARY, cv::THRESH_BINARY_INV, cv::THRESH_TRUNC, cv::THRESH_TOZERO, cv::THRESH_TOZERO_INV) +#define ALL_THRESH_OPS testing::Values(ThreshOp(cv::THRESH_BINARY), ThreshOp(cv::THRESH_BINARY_INV), ThreshOp(cv::THRESH_TRUNC), ThreshOp(cv::THRESH_TOZERO), ThreshOp(cv::THRESH_TOZERO_INV)) CV_ENUM(Interpolation, cv::INTER_NEAREST, cv::INTER_LINEAR, cv::INTER_CUBIC) @@ -194,12 +204,4 @@ CV_FLAGS(DftFlags, cv::DFT_INVERSE, cv::DFT_SCALE, cv::DFT_ROWS, cv::DFT_COMPLEX \ std::make_pair(MatDepth(CV_64F), MatDepth(CV_64F))) -#define WHOLE testing::Values(UseRoi(false)) -#define SUBMAT testing::Values(UseRoi(true)) -#define WHOLE_SUBMAT testing::Values(UseRoi(false), UseRoi(true)) - -#define DIRECT_INVERSE testing::Values(Inverse(false), Inverse(true)) - -#define ALL_CMP_CODES testing::Values(CmpCode(cv::CMP_EQ), CmpCode(cv::CMP_NE), CmpCode(cv::CMP_GT), CmpCode(cv::CMP_GE), CmpCode(cv::CMP_LT), CmpCode(cv::CMP_LE)) - #endif // __OPENCV_TEST_UTILITY_HPP__