From c1c3139368c97f1d31fc323a90b1e99879a79ba4 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Fri, 28 Feb 2014 14:15:56 +0400 Subject: [PATCH 01/10] master-like performance tests --- modules/ocl/perf/perf_arithm.cpp | 291 +++++++++++------- modules/ocl/perf/perf_bgfg.cpp | 24 +- modules/ocl/perf/perf_blend.cpp | 6 +- modules/ocl/perf/perf_brute_force_matcher.cpp | 35 +-- modules/ocl/perf/perf_canny.cpp | 19 +- modules/ocl/perf/perf_color.cpp | 50 +-- modules/ocl/perf/perf_fft.cpp | 17 +- modules/ocl/perf/perf_filters.cpp | 231 ++++++++------ modules/ocl/perf/perf_gemm.cpp | 15 +- modules/ocl/perf/perf_gftt.cpp | 28 +- modules/ocl/perf/perf_haar.cpp | 38 ++- modules/ocl/perf/perf_hog.cpp | 4 +- modules/ocl/perf/perf_imgproc.cpp | 79 +++-- modules/ocl/perf/perf_imgwarp.cpp | 97 +++--- modules/ocl/perf/perf_match_template.cpp | 12 +- modules/ocl/perf/perf_matrix_operation.cpp | 30 +- modules/ocl/perf/perf_moments.cpp | 20 +- modules/ocl/perf/perf_norm.cpp | 27 +- modules/ocl/perf/perf_opticalflow.cpp | 60 +--- modules/ocl/perf/perf_precomp.hpp | 41 ++- modules/ocl/perf/perf_pyramid.cpp | 18 +- modules/ocl/perf/perf_split_merge.cpp | 81 +++-- 22 files changed, 667 insertions(+), 556 deletions(-) diff --git a/modules/ocl/perf/perf_arithm.cpp b/modules/ocl/perf/perf_arithm.cpp index 2699b44a7..7c194ae16 100644 --- a/modules/ocl/perf/perf_arithm.cpp +++ b/modules/ocl/perf/perf_arithm.cpp @@ -54,9 +54,8 @@ using std::tr1::tuple; typedef Size_MatType LUTFixture; -PERF_TEST_P(LUTFixture, LUT, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC3))) +OCL_PERF_TEST_P(LUTFixture, LUT, + ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) { // getting params const Size_MatType_t params = GetParam(); @@ -64,7 +63,7 @@ PERF_TEST_P(LUTFixture, LUT, const int type = get<1>(params); // creating src data - Mat src(srcSize, type), lut(1, 256, CV_8UC1); + Mat src(srcSize, CV_8UC1), lut(1, 256, type); int dstType = CV_MAKETYPE(lut.depth(), src.channels()); Mat dst(srcSize, dstType); @@ -93,16 +92,19 @@ PERF_TEST_P(LUTFixture, LUT, ///////////// Exp //////////////////////// -typedef TestBaseWithParam ExpFixture; +typedef Size_MatType ExpFixture; -PERF_TEST_P(ExpFixture, Exp, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(ExpFixture, Exp, ::testing::Combine( + OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32FC1, CV_32FC4))) { // getting params - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); const double eps = 1e-6; // creating src data - Mat src(srcSize, CV_32FC1), dst(srcSize, CV_32FC1); + Mat src(srcSize, type), dst(srcSize, type); declare.in(src).out(dst); randu(src, 5, 16); @@ -125,18 +127,21 @@ PERF_TEST_P(ExpFixture, Exp, OCL_TYPICAL_MAT_SIZES) SANITY_CHECK(dst, eps, ERROR_RELATIVE); } -///////////// LOG //////////////////////// +///////////// Log //////////////////////// -typedef TestBaseWithParam LogFixture; +typedef Size_MatType LogFixture; -PERF_TEST_P(LogFixture, Log, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(LogFixture, Log, ::testing::Combine( + OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32FC1, CV_32FC4))) { // getting params - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); const double eps = 1e-6; // creating src data - Mat src(srcSize, CV_32F), dst(srcSize, src.type()); + Mat src(srcSize, type), dst(srcSize, type); randu(src, 1, 10); declare.in(src).out(dst); @@ -166,9 +171,8 @@ PERF_TEST_P(LogFixture, Log, OCL_TYPICAL_MAT_SIZES) typedef Size_MatType AddFixture; -PERF_TEST_P(AddFixture, Add, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +OCL_PERF_TEST_P(AddFixture, Add, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134)) { // getting params const Size_MatType_t params = GetParam(); @@ -202,12 +206,51 @@ PERF_TEST_P(AddFixture, Add, OCL_PERF_ELSE } +///////////// Subtract //////////////////////// + +typedef Size_MatType SubtractFixture; + +OCL_PERF_TEST_P(SubtractFixture, Subtract, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134)) +{ + // getting params + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); + + // creating src data + Mat src1(srcSize, type), src2(srcSize, type), dst(srcSize, type); + randu(src1, 0, 1); + randu(src2, 0, 1); + declare.in(src1, src2).out(dst); + + // select implementation + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc1(src1), oclSrc2(src2), oclDst(srcSize, type); + + OCL_TEST_CYCLE() cv::ocl::subtract(oclSrc1, oclSrc2, oclDst); + + oclDst.download(dst); + + SANITY_CHECK(dst); + } + else if (RUN_PLAIN_IMPL) + { + TEST_CYCLE() cv::subtract(src1, src2, dst); + + SANITY_CHECK(dst); + } + else + OCL_PERF_ELSE +} + + ///////////// Mul //////////////////////// typedef Size_MatType MulFixture; -PERF_TEST_P(MulFixture, Mul, ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(MulFixture, Multiply, ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134)) { // getting params const Size_MatType_t params = GetParam(); @@ -245,9 +288,8 @@ PERF_TEST_P(MulFixture, Mul, ::testing::Combine(OCL_TYPICAL_MAT_SIZES, typedef Size_MatType DivFixture; -PERF_TEST_P(DivFixture, Div, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(DivFixture, Divide, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134)) { // getting params const Size_MatType_t params = GetParam(); @@ -260,12 +302,6 @@ PERF_TEST_P(DivFixture, Div, randu(src1, 0, 256); randu(src2, 0, 256); - if ((srcSize == OCL_SIZE_4000 && type == CV_8UC1) || - (srcSize == OCL_SIZE_2000 && type == CV_8UC4)) - declare.time(4.2); - else if (srcSize == OCL_SIZE_4000 && type == CV_8UC4) - declare.time(16.6); - // select implementation if (RUN_OCL_IMPL) { @@ -275,13 +311,13 @@ PERF_TEST_P(DivFixture, Div, oclDst.download(dst); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE); } else if (RUN_PLAIN_IMPL) { TEST_CYCLE() cv::divide(src1, src2, dst); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE); } else OCL_PERF_ELSE @@ -291,9 +327,8 @@ PERF_TEST_P(DivFixture, Div, typedef Size_MatType AbsDiffFixture; -PERF_TEST_P(AbsDiffFixture, Absdiff, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(AbsDiffFixture, Absdiff, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -326,15 +361,18 @@ PERF_TEST_P(AbsDiffFixture, Absdiff, ///////////// CartToPolar //////////////////////// -typedef TestBaseWithParam CartToPolarFixture; +typedef Size_MatType CartToPolarFixture; -PERF_TEST_P(CartToPolarFixture, CartToPolar, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(CartToPolarFixture, CartToPolar, ::testing::Combine( + OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32FC1, CV_32FC4))) { - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); const double eps = 8e-3; - Mat src1(srcSize, CV_32FC1), src2(srcSize, CV_32FC1), - dst1(srcSize, CV_32FC1), dst2(srcSize, CV_32FC1); + Mat src1(srcSize, type), src2(srcSize, type), + dst1(srcSize, type), dst2(srcSize, type); declare.in(src1, src2).out(dst1, dst2); randu(src1, 0, 256); randu(src2, 0, 256); @@ -368,14 +406,17 @@ PERF_TEST_P(CartToPolarFixture, CartToPolar, OCL_TYPICAL_MAT_SIZES) ///////////// PolarToCart //////////////////////// -typedef TestBaseWithParam PolarToCartFixture; +typedef Size_MatType PolarToCartFixture; -PERF_TEST_P(PolarToCartFixture, PolarToCart, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(PolarToCartFixture, PolarToCart, ::testing::Combine( + OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32FC1, CV_32FC4))) { - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); - Mat src1(srcSize, CV_32FC1), src2(srcSize, CV_32FC1), - dst1(srcSize, CV_32FC1), dst2(srcSize, CV_32FC1); + Mat src1(srcSize, type), src2(srcSize, type), + dst1(srcSize, type), dst2(srcSize, type); declare.in(src1, src2).out(dst1, dst2); randu(src1, 0, 256); randu(src2, 0, 256); @@ -409,14 +450,17 @@ PERF_TEST_P(PolarToCartFixture, PolarToCart, OCL_TYPICAL_MAT_SIZES) ///////////// Magnitude //////////////////////// -typedef TestBaseWithParam MagnitudeFixture; +typedef Size_MatType MagnitudeFixture; -PERF_TEST_P(MagnitudeFixture, Magnitude, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(MagnitudeFixture, Magnitude, ::testing::Combine( + OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32FC1, CV_32FC4))) { - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); - Mat src1(srcSize, CV_32FC1), src2(srcSize, CV_32FC1), - dst(srcSize, CV_32FC1); + Mat src1(srcSize, type), src2(srcSize, type), + dst(srcSize, type); randu(src1, 0, 1); randu(src2, 0, 1); declare.in(src1, src2).out(dst); @@ -446,9 +490,8 @@ PERF_TEST_P(MagnitudeFixture, Magnitude, OCL_TYPICAL_MAT_SIZES) typedef Size_MatType TransposeFixture; -PERF_TEST_P(TransposeFixture, Transpose, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(TransposeFixture, Transpose, ::testing::Combine( + OCL_TEST_SIZES, OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -479,15 +522,24 @@ PERF_TEST_P(TransposeFixture, Transpose, ///////////// Flip //////////////////////// -typedef Size_MatType FlipFixture; - -PERF_TEST_P(FlipFixture, Flip, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +enum { - const Size_MatType_t params = GetParam(); + FLIP_BOTH = 0, FLIP_ROWS, FLIP_COLS +}; + +CV_ENUM(FlipType, FLIP_BOTH, FLIP_ROWS, FLIP_COLS) + +typedef std::tr1::tuple FlipParams; +typedef TestBaseWithParam FlipFixture; + +OCL_PERF_TEST_P(FlipFixture, Flip, + ::testing::Combine(OCL_TEST_SIZES, + OCL_TEST_TYPES, FlipType::all())) +{ + const FlipParams params = GetParam(); const Size srcSize = get<0>(params); const int type = get<1>(params); + const int flipType = get<2>(params); Mat src(srcSize, type), dst(srcSize, type); declare.in(src, WARMUP_RNG).out(dst); @@ -496,7 +548,7 @@ PERF_TEST_P(FlipFixture, Flip, { ocl::oclMat oclSrc(src), oclDst(srcSize, type); - OCL_TEST_CYCLE() cv::ocl::flip(oclSrc, oclDst, 0); + OCL_TEST_CYCLE() cv::ocl::flip(oclSrc, oclDst, flipType - 1); oclDst.download(dst); @@ -504,7 +556,7 @@ PERF_TEST_P(FlipFixture, Flip, } else if (RUN_PLAIN_IMPL) { - TEST_CYCLE() cv::flip(src, dst, 0); + TEST_CYCLE() cv::flip(src, dst, flipType - 1); SANITY_CHECK(dst); } @@ -512,11 +564,11 @@ PERF_TEST_P(FlipFixture, Flip, OCL_PERF_ELSE } -///////////// minMax //////////////////////// +///////////// MinMax //////////////////////// -typedef Size_MatType minMaxFixture; +typedef Size_MatType MinMaxFixture; -PERF_TEST_P(minMaxFixture, minMax, +PERF_TEST_P(MinMaxFixture, MinMax, ::testing::Combine(OCL_TYPICAL_MAT_SIZES, OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) { @@ -554,13 +606,12 @@ PERF_TEST_P(minMaxFixture, minMax, OCL_PERF_ELSE } -///////////// minMaxLoc //////////////////////// +///////////// MinMaxLoc //////////////////////// -typedef Size_MatType minMaxLocFixture; +typedef Size_MatType MinMaxLocFixture; -PERF_TEST_P(minMaxLocFixture, minMaxLoc, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +OCL_PERF_TEST_P(MinMaxLocFixture, MinMaxLoc, + ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -599,9 +650,9 @@ PERF_TEST_P(minMaxLocFixture, minMaxLoc, typedef Size_MatType SumFixture; -PERF_TEST_P(SumFixture, Sum, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32SC1))) +OCL_PERF_TEST_P(SumFixture, Sum, + ::testing::Combine(OCL_TEST_SIZES, + OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -618,13 +669,13 @@ PERF_TEST_P(SumFixture, Sum, OCL_TEST_CYCLE() result = cv::ocl::sum(oclSrc); - SANITY_CHECK(result); + SANITY_CHECK(result, 1e-6, ERROR_RELATIVE); } else if (RUN_PLAIN_IMPL) { TEST_CYCLE() result = cv::sum(src); - SANITY_CHECK(result); + SANITY_CHECK(result, 1e-6, ERROR_RELATIVE); } else OCL_PERF_ELSE @@ -632,10 +683,10 @@ PERF_TEST_P(SumFixture, Sum, ///////////// countNonZero //////////////////////// -typedef Size_MatType countNonZeroFixture; +typedef Size_MatType CountNonZeroFixture; -PERF_TEST_P(countNonZeroFixture, countNonZero, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, +OCL_PERF_TEST_P(CountNonZeroFixture, CountNonZero, + ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) { const Size_MatType_t params = GetParam(); @@ -667,14 +718,16 @@ PERF_TEST_P(countNonZeroFixture, countNonZero, ///////////// Phase //////////////////////// -typedef TestBaseWithParam PhaseFixture; +typedef Size_MatType PhaseFixture; -PERF_TEST_P(PhaseFixture, Phase, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(PhaseFixture, Phase, ::testing::Combine( + OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32FC1, CV_32FC4))) { - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); - Mat src1(srcSize, CV_32FC1), src2(srcSize, CV_32FC1), - dst(srcSize, CV_32FC1); + Mat src1(srcSize, type), src2(srcSize, type), dst(srcSize, type); declare.in(src1, src2).out(dst); randu(src1, 0, 256); randu(src2, 0, 256); @@ -704,9 +757,8 @@ PERF_TEST_P(PhaseFixture, Phase, OCL_TYPICAL_MAT_SIZES) typedef Size_MatType BitwiseAndFixture; -PERF_TEST_P(BitwiseAndFixture, bitwise_and, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32SC1))) +OCL_PERF_TEST_P(BitwiseAndFixture, Bitwise_and, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -741,9 +793,8 @@ PERF_TEST_P(BitwiseAndFixture, bitwise_and, typedef Size_MatType BitwiseXorFixture; -PERF_TEST_P(BitwiseXorFixture, bitwise_xor, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32SC1))) +OCL_PERF_TEST_P(BitwiseXorFixture, Bitwise_xor, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -778,9 +829,8 @@ PERF_TEST_P(BitwiseXorFixture, bitwise_xor, typedef Size_MatType BitwiseOrFixture; -PERF_TEST_P(BitwiseOrFixture, bitwise_or, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32SC1))) +OCL_PERF_TEST_P(BitwiseOrFixture, Bitwise_or, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -798,26 +848,25 @@ PERF_TEST_P(BitwiseOrFixture, bitwise_or, OCL_TEST_CYCLE() cv::ocl::bitwise_or(oclSrc1, oclSrc2, oclDst); oclDst.download(dst); - - SANITY_CHECK(dst); } else if (RUN_PLAIN_IMPL) { TEST_CYCLE() cv::bitwise_or(src1, src2, dst); - - SANITY_CHECK(dst); } else OCL_PERF_ELSE + + if (CV_MAT_DEPTH(type) >= CV_32F) + cv::patchNaNs(dst, 17); + SANITY_CHECK(dst); } ///////////// bitwise_not//////////////////////// typedef Size_MatType BitwiseNotFixture; -PERF_TEST_P(BitwiseAndFixture, bitwise_not, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32SC1))) +OCL_PERF_TEST_P(BitwiseNotFixture, Bitwise_not, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -848,15 +897,19 @@ PERF_TEST_P(BitwiseAndFixture, bitwise_not, ///////////// compare//////////////////////// -typedef Size_MatType CompareFixture; +CV_ENUM(CmpCode, CMP_LT, CMP_LE, CMP_EQ, CMP_NE, CMP_GE, CMP_GT) -PERF_TEST_P(CompareFixture, compare, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +typedef std::tr1::tuple CompareParams; +typedef TestBaseWithParam CompareFixture; + +OCL_PERF_TEST_P(CompareFixture, Compare, + ::testing::Combine(OCL_TEST_SIZES, + OCL_PERF_ENUM(CV_8UC1, CV_32FC1), CmpCode::all())) { - const Size_MatType_t params = GetParam(); + const CompareParams params = GetParam(); const Size srcSize = get<0>(params); const int type = get<1>(params); + const int cmpCode = get<2>(params); Mat src1(srcSize, type), src2(srcSize, type), dst(srcSize, CV_8UC1); declare.in(src1, src2, WARMUP_RNG).out(dst); @@ -865,7 +918,7 @@ PERF_TEST_P(CompareFixture, compare, { ocl::oclMat oclSrc1(src1), oclSrc2(src2), oclDst(srcSize, CV_8UC1); - OCL_TEST_CYCLE() cv::ocl::compare(oclSrc1, oclSrc2, oclDst, CMP_EQ); + OCL_TEST_CYCLE() cv::ocl::compare(oclSrc1, oclSrc2, oclDst, cmpCode); oclDst.download(dst); @@ -873,7 +926,7 @@ PERF_TEST_P(CompareFixture, compare, } else if (RUN_PLAIN_IMPL) { - TEST_CYCLE() cv::compare(src1, src2, dst, CMP_EQ); + TEST_CYCLE() cv::compare(src1, src2, dst, cmpCode); SANITY_CHECK(dst); } @@ -883,14 +936,17 @@ PERF_TEST_P(CompareFixture, compare, ///////////// pow //////////////////////// -typedef TestBaseWithParam PowFixture; +typedef Size_MatType PowFixture; -PERF_TEST_P(PowFixture, pow, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(PowFixture, Pow, ::testing::Combine( + OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32FC1, CV_32FC4))) { - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); const double eps = 1e-6; - Mat src(srcSize, CV_32F), dst(srcSize, CV_32F); + Mat src(srcSize, type), dst(srcSize, type); declare.in(src, WARMUP_RNG).out(dst); if (RUN_OCL_IMPL) @@ -915,9 +971,8 @@ PERF_TEST_P(PowFixture, pow, OCL_TYPICAL_MAT_SIZES) typedef Size_MatType AddWeightedFixture; -PERF_TEST_P(AddWeightedFixture, AddWeighted, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +OCL_PERF_TEST_P(AddWeightedFixture, AddWeighted, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -951,9 +1006,8 @@ PERF_TEST_P(AddWeightedFixture, AddWeighted, typedef Size_MatType MinFixture; -PERF_TEST_P(MinFixture, Min, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +OCL_PERF_TEST_P(MinFixture, Min, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -986,9 +1040,8 @@ PERF_TEST_P(MinFixture, Min, typedef Size_MatType MaxFixture; -PERF_TEST_P(MaxFixture, Max, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +OCL_PERF_TEST_P(MaxFixture, Max, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -1017,7 +1070,7 @@ PERF_TEST_P(MaxFixture, Max, OCL_PERF_ELSE } -///////////// Max //////////////////////// +///////////// Abs //////////////////////// typedef Size_MatType AbsFixture; @@ -1056,9 +1109,9 @@ PERF_TEST_P(AbsFixture, Abs, typedef Size_MatType RepeatFixture; -PERF_TEST_P(RepeatFixture, Repeat, - ::testing::Combine(::testing::Values(OCL_SIZE_1000, OCL_SIZE_2000), - OCL_PERF_ENUM(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4))) +OCL_PERF_TEST_P(RepeatFixture, Repeat, + ::testing::Combine(OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3), + OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); diff --git a/modules/ocl/perf/perf_bgfg.cpp b/modules/ocl/perf/perf_bgfg.cpp index b361a9232..c2a1941e9 100644 --- a/modules/ocl/perf/perf_bgfg.cpp +++ b/modules/ocl/perf/perf_bgfg.cpp @@ -165,11 +165,11 @@ PERF_TEST_P(VideoMOGFixture, MOG, ///////////// MOG2 //////////////////////// typedef tuple VideoMOG2ParamType; -typedef TestBaseWithParam VideoMOG2Fixture; +typedef TestBaseWithParam MOG2_Apply; -PERF_TEST_P(VideoMOG2Fixture, DISABLED_MOG2, // TODO Disabled: random hungs on buildslave - ::testing::Combine(::testing::Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"), - ::testing::Values(1, 3))) +OCL_PERF_TEST_P(MOG2_Apply, Mog2, + testing::Combine(testing::Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"), + testing::Values(1, 3))) { VideoMOG2ParamType params = GetParam(); @@ -195,9 +195,7 @@ PERF_TEST_P(VideoMOG2Fixture, DISABLED_MOG2, // TODO Disabled: random hungs on b foreground.release(); for (int i = 0; i < nFrame; i++) - { mog2(frame_buffer[i], foreground); - } } SANITY_CHECK(foreground); } @@ -210,9 +208,7 @@ PERF_TEST_P(VideoMOG2Fixture, DISABLED_MOG2, // TODO Disabled: random hungs on b cv::ocl::MOG2 d_mog2; foreground_d.release(); for (int i = 0; i < nFrame; i++) - { d_mog2(frame_buffer_ocl[i], foreground_d); - } } foreground_d.download(foreground); SANITY_CHECK(foreground); @@ -223,11 +219,11 @@ PERF_TEST_P(VideoMOG2Fixture, DISABLED_MOG2, // TODO Disabled: random hungs on b ///////////// MOG2_GetBackgroundImage ////////////////// -typedef TestBaseWithParam Video_MOG2GetBackgroundImage; +typedef TestBaseWithParam MOG2_GetBackgroundImage; -PERF_TEST_P(Video_MOG2GetBackgroundImage, MOG2, - ::testing::Combine(::testing::Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"), - ::testing::Values(3))) +OCL_PERF_TEST_P(MOG2_GetBackgroundImage, Mog2, + testing::Combine(testing::Values("gpu/video/768x576.avi", "gpu/video/1920x1080.avi"), + testing::Values(3))) { VideoMOG2ParamType params = GetParam(); @@ -248,7 +244,7 @@ PERF_TEST_P(Video_MOG2GetBackgroundImage, MOG2, cv::ocl::oclMat foreground_d; cv::ocl::oclMat background_d; - if(RUN_PLAIN_IMPL) + if (RUN_PLAIN_IMPL) { TEST_CYCLE() { @@ -264,7 +260,7 @@ PERF_TEST_P(Video_MOG2GetBackgroundImage, MOG2, } SANITY_CHECK(background); } - else if(RUN_OCL_IMPL) + else if (RUN_OCL_IMPL) { prepareData(frame_buffer, frame_buffer_ocl); CV_Assert((int)(frame_buffer_ocl.size()) == nFrame); diff --git a/modules/ocl/perf/perf_blend.cpp b/modules/ocl/perf/perf_blend.cpp index 6f611bbc3..5471dee2e 100644 --- a/modules/ocl/perf/perf_blend.cpp +++ b/modules/ocl/perf/perf_blend.cpp @@ -88,10 +88,10 @@ typedef void (*blendFunction)(const Mat &img1, const Mat &img2, const Mat &weights1, const Mat &weights2, Mat &result_gold); -typedef Size_MatType blendLinearFixture; +typedef Size_MatType BlendLinearFixture; -PERF_TEST_P(blendLinearFixture, blendLinear, ::testing::Combine( - OCL_TYPICAL_MAT_SIZES, testing::Values(CV_8UC1, CV_8UC3, CV_32FC1))) +OCL_PERF_TEST_P(BlendLinearFixture, BlendLinear, + ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32FC1, CV_32FC4))) { Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); diff --git a/modules/ocl/perf/perf_brute_force_matcher.cpp b/modules/ocl/perf/perf_brute_force_matcher.cpp index d124428e9..54e829f4b 100644 --- a/modules/ocl/perf/perf_brute_force_matcher.cpp +++ b/modules/ocl/perf/perf_brute_force_matcher.cpp @@ -47,20 +47,17 @@ using namespace perf; -#define OCL_BFMATCHER_TYPICAL_MAT_SIZES ::testing::Values(cv::Size(128, 500), cv::Size(128, 1000), cv::Size(128, 2000)) - //////////////////// BruteForceMatch ///////////////// typedef TestBaseWithParam BruteForceMatcherFixture; -PERF_TEST_P(BruteForceMatcherFixture, match, - OCL_BFMATCHER_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(BruteForceMatcherFixture, Match, OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3)) { const Size srcSize = GetParam(); vector matches; - Mat query(srcSize, CV_32F), train(srcSize, CV_32F); - declare.in(query, train).time(srcSize.height == 2000 ? 9 : 4 ); + Mat query(srcSize, CV_32FC1), train(srcSize, CV_32FC1); + declare.in(query, train); randu(query, 0.0f, 1.0f); randu(train, 0.0f, 1.0f); @@ -75,12 +72,9 @@ PERF_TEST_P(BruteForceMatcherFixture, match, { ocl::BruteForceMatcher_OCL_base oclMatcher(ocl::BruteForceMatcher_OCL_base::L2Dist); ocl::oclMat oclQuery(query), oclTrain(train); - ocl::oclMat oclTrainIdx, oclDistance; OCL_TEST_CYCLE() - oclMatcher.matchSingle(oclQuery, oclTrain, oclTrainIdx, oclDistance); - - oclMatcher.matchDownload(oclTrainIdx, oclDistance, matches); + oclMatcher.match(oclQuery, oclTrain, matches); SANITY_CHECK_MATCHES(matches, 1e-5); } @@ -88,8 +82,7 @@ PERF_TEST_P(BruteForceMatcherFixture, match, OCL_PERF_ELSE } -PERF_TEST_P(BruteForceMatcherFixture, knnMatch, - OCL_BFMATCHER_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(BruteForceMatcherFixture, KnnMatch, OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3)) { const Size srcSize = GetParam(); @@ -99,8 +92,6 @@ PERF_TEST_P(BruteForceMatcherFixture, knnMatch, randu(train, 0.0f, 1.0f); declare.in(query, train); - if (srcSize.height == 2000) - declare.time(9); if (RUN_PLAIN_IMPL) { @@ -115,10 +106,10 @@ PERF_TEST_P(BruteForceMatcherFixture, knnMatch, { ocl::BruteForceMatcher_OCL_base oclMatcher(ocl::BruteForceMatcher_OCL_base::L2Dist); ocl::oclMat oclQuery(query), oclTrain(train); - ocl::oclMat oclTrainIdx, oclDistance, oclAllDist; + ocl::oclMat oclTrainIdx, oclDistance; OCL_TEST_CYCLE() - oclMatcher.knnMatchSingle(oclQuery, oclTrain, oclTrainIdx, oclDistance, oclAllDist, 2); + oclMatcher.knnMatch(oclQuery, oclTrain, matches, 2); oclMatcher.knnMatchDownload(oclTrainIdx, oclDistance, matches); @@ -130,22 +121,18 @@ PERF_TEST_P(BruteForceMatcherFixture, knnMatch, OCL_PERF_ELSE } -PERF_TEST_P(BruteForceMatcherFixture, radiusMatch, - OCL_BFMATCHER_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(BruteForceMatcherFixture, RadiusMatch, OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3)) { const Size srcSize = GetParam(); const float max_distance = 2.0f; vector > matches(2); - Mat query(srcSize, CV_32F), train(srcSize, CV_32F); + Mat query(srcSize, CV_32FC1), train(srcSize, CV_32FC1); declare.in(query, train); randu(query, 0.0f, 1.0f); randu(train, 0.0f, 1.0f); - if (srcSize.height == 2000) - declare.time(9.15); - if (RUN_PLAIN_IMPL) { cv::BFMatcher matcher(NORM_L2); @@ -162,7 +149,7 @@ PERF_TEST_P(BruteForceMatcherFixture, radiusMatch, ocl::oclMat oclTrainIdx, oclDistance, oclNMatches; OCL_TEST_CYCLE() - oclMatcher.radiusMatchSingle(oclQuery, oclTrain, oclTrainIdx, oclDistance, oclNMatches, max_distance); + oclMatcher.radiusMatch(oclQuery, oclTrain, matches, max_distance); oclMatcher.radiusMatchDownload(oclTrainIdx, oclDistance, oclNMatches, matches); @@ -173,5 +160,3 @@ PERF_TEST_P(BruteForceMatcherFixture, radiusMatch, else OCL_PERF_ELSE } - -#undef OCL_BFMATCHER_TYPICAL_MAT_SIZES diff --git a/modules/ocl/perf/perf_canny.cpp b/modules/ocl/perf/perf_canny.cpp index 33723daa3..96c444151 100644 --- a/modules/ocl/perf/perf_canny.cpp +++ b/modules/ocl/perf/perf_canny.cpp @@ -46,11 +46,21 @@ #include "perf_precomp.hpp" using namespace perf; +using std::tr1::tuple; +using std::tr1::get; ///////////// Canny //////////////////////// -PERF_TEST(CannyFixture, Canny) +typedef tuple CannyParams; +typedef TestBaseWithParam CannyFixture; + +OCL_PERF_TEST_P(CannyFixture, Canny, + ::testing::Combine(OCL_PERF_ENUM(3, 5), testing::Bool())) { + const CannyParams params = GetParam(); + int apertureSize = get<0>(params); + bool L2Grad = get<1>(params); + Mat img = imread(getDataPath("gpu/stereobm/aloe-L.png"), cv::IMREAD_GRAYSCALE), edges(img.size(), CV_8UC1); ASSERT_TRUE(!img.empty()) << "can't open aloe-L.png"; @@ -61,16 +71,15 @@ PERF_TEST(CannyFixture, Canny) { ocl::oclMat oclImg(img), oclEdges(img.size(), CV_8UC1); - OCL_TEST_CYCLE() ocl::Canny(oclImg, oclEdges, 50.0, 100.0); + OCL_TEST_CYCLE() ocl::Canny(oclImg, oclEdges, 50.0, 100.0, apertureSize, L2Grad); oclEdges.download(edges); } else if (RUN_PLAIN_IMPL) { - TEST_CYCLE() Canny(img, edges, 50.0, 100.0); + TEST_CYCLE() Canny(img, edges, 50.0, 100.0, apertureSize, L2Grad); } else OCL_PERF_ELSE - int value = 0; - SANITY_CHECK(value); + SANITY_CHECK_NOTHING(); } diff --git a/modules/ocl/perf/perf_color.cpp b/modules/ocl/perf/perf_color.cpp index 843331518..0797a87fd 100644 --- a/modules/ocl/perf/perf_color.cpp +++ b/modules/ocl/perf/perf_color.cpp @@ -52,36 +52,36 @@ using std::tr1::make_tuple; ///////////// cvtColor//////////////////////// -CV_ENUM(ConversionTypes, CV_RGB2GRAY, CV_RGB2BGR, CV_RGB2YUV, CV_YUV2RGB, CV_RGB2YCrCb, - CV_YCrCb2RGB, CV_RGB2XYZ, CV_XYZ2RGB, CV_RGB2HSV, CV_HSV2RGB, CV_RGB2HLS, - CV_HLS2RGB, CV_BGR5652BGR, CV_BGR2BGR565, CV_RGBA2mRGBA, CV_mRGBA2RGBA, CV_YUV2RGB_NV12) +CV_ENUM(ConversionTypes, COLOR_RGB2GRAY, COLOR_RGB2BGR, COLOR_RGB2YUV, COLOR_YUV2RGB, COLOR_RGB2YCrCb, + COLOR_YCrCb2RGB, COLOR_RGB2XYZ, COLOR_XYZ2RGB, COLOR_RGB2HSV, COLOR_HSV2RGB, COLOR_RGB2HLS, + COLOR_HLS2RGB, COLOR_BGR5652BGR, COLOR_BGR2BGR565, COLOR_RGBA2mRGBA, COLOR_mRGBA2RGBA, COLOR_YUV2RGB_NV12) -typedef tuple > cvtColorParams; -typedef TestBaseWithParam cvtColorFixture; +typedef tuple > CvtColorParams; +typedef TestBaseWithParam CvtColorFixture; -PERF_TEST_P(cvtColorFixture, cvtColor, testing::Combine( - testing::Values(Size(1000, 1002), Size(2000, 2004), Size(4000, 4008)), +OCL_PERF_TEST_P(CvtColorFixture, CvtColor, testing::Combine( + OCL_TEST_SIZES, testing::Values( - make_tuple(ConversionTypes(CV_RGB2GRAY), 3, 1), - make_tuple(ConversionTypes(CV_RGB2BGR), 3, 3), - make_tuple(ConversionTypes(CV_RGB2YUV), 3, 3), - make_tuple(ConversionTypes(CV_YUV2RGB), 3, 3), - make_tuple(ConversionTypes(CV_RGB2YCrCb), 3, 3), - make_tuple(ConversionTypes(CV_YCrCb2RGB), 3, 3), - make_tuple(ConversionTypes(CV_RGB2XYZ), 3, 3), - make_tuple(ConversionTypes(CV_XYZ2RGB), 3, 3), - make_tuple(ConversionTypes(CV_RGB2HSV), 3, 3), - make_tuple(ConversionTypes(CV_HSV2RGB), 3, 3), - make_tuple(ConversionTypes(CV_RGB2HLS), 3, 3), - make_tuple(ConversionTypes(CV_HLS2RGB), 3, 3), - make_tuple(ConversionTypes(CV_BGR5652BGR), 2, 3), - make_tuple(ConversionTypes(CV_BGR2BGR565), 3, 2), - make_tuple(ConversionTypes(CV_RGBA2mRGBA), 4, 4), - make_tuple(ConversionTypes(CV_mRGBA2RGBA), 4, 4), - make_tuple(ConversionTypes(CV_YUV2RGB_NV12), 1, 3) + make_tuple(ConversionTypes(COLOR_RGB2GRAY), 3, 1), + make_tuple(ConversionTypes(COLOR_RGB2BGR), 3, 3), + make_tuple(ConversionTypes(COLOR_RGB2YUV), 3, 3), + make_tuple(ConversionTypes(COLOR_YUV2RGB), 3, 3), + make_tuple(ConversionTypes(COLOR_RGB2YCrCb), 3, 3), + make_tuple(ConversionTypes(COLOR_YCrCb2RGB), 3, 3), + make_tuple(ConversionTypes(COLOR_RGB2XYZ), 3, 3), + make_tuple(ConversionTypes(COLOR_XYZ2RGB), 3, 3), + make_tuple(ConversionTypes(COLOR_RGB2HSV), 3, 3), + make_tuple(ConversionTypes(COLOR_HSV2RGB), 3, 3), + make_tuple(ConversionTypes(COLOR_RGB2HLS), 3, 3), + make_tuple(ConversionTypes(COLOR_HLS2RGB), 3, 3), + make_tuple(ConversionTypes(COLOR_BGR5652BGR), 2, 3), + make_tuple(ConversionTypes(COLOR_BGR2BGR565), 3, 2), + make_tuple(ConversionTypes(COLOR_RGBA2mRGBA), 4, 4), + make_tuple(ConversionTypes(COLOR_mRGBA2RGBA), 4, 4), + make_tuple(ConversionTypes(COLOR_YUV2RGB_NV12), 1, 3) ))) { - cvtColorParams params = GetParam(); + CvtColorParams params = GetParam(); const Size srcSize = get<0>(params); const tuple conversionParams = get<1>(params); const int code = get<0>(conversionParams), scn = get<1>(conversionParams), diff --git a/modules/ocl/perf/perf_fft.cpp b/modules/ocl/perf/perf_fft.cpp index 49da65936..29e042ccd 100644 --- a/modules/ocl/perf/perf_fft.cpp +++ b/modules/ocl/perf/perf_fft.cpp @@ -47,6 +47,8 @@ #include "perf_precomp.hpp" using namespace perf; +using std::tr1::tuple; +using std::tr1::get; ///////////// dft //////////////////////// @@ -54,9 +56,16 @@ typedef TestBaseWithParam dftFixture; #ifdef HAVE_CLAMDFFT -PERF_TEST_P(dftFixture, dft, OCL_TYPICAL_MAT_SIZES) +typedef tuple DftParams; +typedef TestBaseWithParam DftFixture; + +OCL_PERF_TEST_P(DftFixture, Dft, ::testing::Combine(testing::Values(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3), + ::testing::Values((int)DFT_ROWS, (int)DFT_SCALE, (int)DFT_INVERSE, + (int)DFT_INVERSE | DFT_SCALE, (int)DFT_ROWS | DFT_INVERSE))) { - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int flags = get<1>(params); Mat src(srcSize, CV_32FC2), dst; randu(src, 0.0f, 1.0f); @@ -69,7 +78,7 @@ PERF_TEST_P(dftFixture, dft, OCL_TYPICAL_MAT_SIZES) { ocl::oclMat oclSrc(src), oclDst; - OCL_TEST_CYCLE() cv::ocl::dft(oclSrc, oclDst); + OCL_TEST_CYCLE() cv::ocl::dft(oclSrc, oclDst, Size(), flags | DFT_COMPLEX_OUTPUT); oclDst.download(dst); @@ -77,7 +86,7 @@ PERF_TEST_P(dftFixture, dft, OCL_TYPICAL_MAT_SIZES) } else if (RUN_PLAIN_IMPL) { - TEST_CYCLE() cv::dft(src, dst); + TEST_CYCLE() cv::dft(src, dst, flags | DFT_COMPLEX_OUTPUT); SANITY_CHECK(dst); } diff --git a/modules/ocl/perf/perf_filters.cpp b/modules/ocl/perf/perf_filters.cpp index c625caa47..70e51b466 100644 --- a/modules/ocl/perf/perf_filters.cpp +++ b/modules/ocl/perf/perf_filters.cpp @@ -49,31 +49,30 @@ using namespace perf; using std::tr1::get; using std::tr1::tuple; +typedef tuple FilterParams; +typedef TestBaseWithParam FilterFixture; + ///////////// Blur//////////////////////// -typedef Size_MatType BlurFixture; +typedef FilterFixture BlurFixture; -PERF_TEST_P(BlurFixture, Blur, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(BlurFixture, Blur, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5))) { - const Size_MatType_t params = GetParam(); - const Size srcSize = get<0>(params), ksize(3, 3); - const int type = get<1>(params), bordertype = BORDER_CONSTANT; + const FilterParams params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params), ksize = get<2>(params), bordertype = BORDER_CONSTANT; checkDeviceMaxMemoryAllocSize(srcSize, type); Mat src(srcSize, type), dst(srcSize, type); declare.in(src, WARMUP_RNG).out(dst); - if (srcSize == OCL_SIZE_4000 && type == CV_8UC4) - declare.time(5); - if (RUN_OCL_IMPL) { ocl::oclMat oclSrc(src), oclDst(srcSize, type); - OCL_TEST_CYCLE() cv::ocl::blur(oclSrc, oclDst, ksize, Point(-1, -1), bordertype); + OCL_TEST_CYCLE() cv::ocl::blur(oclSrc, oclDst, Size(ksize, ksize), Point(-1, -1), bordertype); oclDst.download(dst); @@ -81,7 +80,7 @@ PERF_TEST_P(BlurFixture, Blur, } else if (RUN_PLAIN_IMPL) { - TEST_CYCLE() cv::blur(src, dst, ksize, Point(-1, -1), bordertype); + TEST_CYCLE() cv::blur(src, dst, Size(ksize, ksize), Point(-1, -1), bordertype); SANITY_CHECK(dst, 1 + DBL_EPSILON); } @@ -91,24 +90,20 @@ PERF_TEST_P(BlurFixture, Blur, ///////////// Laplacian//////////////////////// -typedef Size_MatType LaplacianFixture; +typedef FilterFixture LaplacianFixture; -PERF_TEST_P(LaplacianFixture, Laplacian, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(LaplacianFixture, Laplacian, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(1, 3))) { - const Size_MatType_t params = GetParam(); + const FilterParams params = GetParam(); const Size srcSize = get<0>(params); - const int type = get<1>(params), ksize = 3; + const int type = get<1>(params), ksize = get<2>(params); checkDeviceMaxMemoryAllocSize(srcSize, type); Mat src(srcSize, type), dst(srcSize, type); declare.in(src, WARMUP_RNG).out(dst); - if (srcSize == OCL_SIZE_4000 && type == CV_8UC4) - declare.time(6); - if (RUN_OCL_IMPL) { ocl::oclMat oclSrc(src), oclDst(srcSize, type); @@ -117,13 +112,13 @@ PERF_TEST_P(LaplacianFixture, Laplacian, oclDst.download(dst); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 5e-3); } else if (RUN_PLAIN_IMPL) { TEST_CYCLE() cv::Laplacian(src, dst, -1, ksize, 1); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 5e-3); } else OCL_PERF_ELSE @@ -131,15 +126,14 @@ PERF_TEST_P(LaplacianFixture, Laplacian, ///////////// Erode //////////////////// -typedef Size_MatType ErodeFixture; +typedef FilterFixture ErodeFixture; -PERF_TEST_P(ErodeFixture, Erode, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4))) +OCL_PERF_TEST_P(ErodeFixture, Erode, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5))) { - const Size_MatType_t params = GetParam(); + const FilterParams params = GetParam(); const Size srcSize = get<0>(params); - const int type = get<1>(params), ksize = 3; + const int type = get<1>(params), ksize = get<2>(params); const Mat ker = getStructuringElement(MORPH_RECT, Size(ksize, ksize)); checkDeviceMaxMemoryAllocSize(srcSize, type); @@ -147,9 +141,6 @@ PERF_TEST_P(ErodeFixture, Erode, Mat src(srcSize, type), dst(srcSize, type); declare.in(src, WARMUP_RNG).out(dst).in(ker); - if (srcSize == OCL_SIZE_4000 && type == CV_8UC4) - declare.time(5); - if (RUN_OCL_IMPL) { ocl::oclMat oclSrc(src), oclDst(srcSize, type), oclKer(ker); @@ -170,13 +161,89 @@ PERF_TEST_P(ErodeFixture, Erode, OCL_PERF_ELSE } +///////////// Dilate //////////////////// + +typedef FilterFixture DilateFixture; + +OCL_PERF_TEST_P(DilateFixture, Dilate, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5))) +{ + const FilterParams params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params), ksize = get<2>(params); + const Mat ker = getStructuringElement(MORPH_RECT, Size(ksize, ksize)); + + checkDeviceMaxMemoryAllocSize(srcSize, type); + + Mat src(srcSize, type), dst(srcSize, type); + declare.in(src, WARMUP_RNG).out(dst).in(ker); + + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc(src), oclDst(srcSize, type), oclKer(ker); + + OCL_TEST_CYCLE() cv::ocl::dilate(oclSrc, oclDst, oclKer); + + oclDst.download(dst); + + SANITY_CHECK(dst); + } + else if (RUN_PLAIN_IMPL) + { + TEST_CYCLE() cv::dilate(src, dst, ker); + + SANITY_CHECK(dst); + } + else + OCL_PERF_ELSE +} + +///////////// MorphologyEx //////////////////// + +CV_ENUM(MorphOp, MORPH_OPEN, MORPH_CLOSE, MORPH_GRADIENT, MORPH_TOPHAT, MORPH_BLACKHAT) + +typedef tuple MorphologyExParams; +typedef TestBaseWithParam MorphologyExFixture; + +OCL_PERF_TEST_P(MorphologyExFixture, MorphologyEx, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, MorphOp::all(), OCL_PERF_ENUM(3, 5))) +{ + const MorphologyExParams params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params), op = get<2>(params), ksize = get<3>(params); + const Mat ker = getStructuringElement(MORPH_RECT, Size(ksize, ksize)); + + checkDeviceMaxMemoryAllocSize(srcSize, type); + + Mat src(srcSize, type), dst(srcSize, type); + declare.in(src, WARMUP_RNG).out(dst).in(ker); + + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc(src), oclDst(srcSize, type), oclKer(ker); + + OCL_TEST_CYCLE() cv::ocl::morphologyEx(oclSrc, oclDst, op, oclKer); + + oclDst.download(dst); + + SANITY_CHECK(dst); + } + else if (RUN_PLAIN_IMPL) + { + TEST_CYCLE() cv::morphologyEx(src, dst, op, ker); + + SANITY_CHECK(dst); + } + else + OCL_PERF_ELSE +} + ///////////// Sobel //////////////////////// typedef Size_MatType SobelFixture; -PERF_TEST_P(SobelFixture, Sobel, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(SobelFixture, Sobel, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -187,12 +254,6 @@ PERF_TEST_P(SobelFixture, Sobel, Mat src(srcSize, type), dst(srcSize, type); declare.in(src, WARMUP_RNG).out(dst); - if ((srcSize == OCL_SIZE_2000 && type == CV_8UC4) || - (srcSize == OCL_SIZE_4000 && type == CV_8UC1)) - declare.time(5.5); - else if (srcSize == OCL_SIZE_4000 && type == CV_8UC4) - declare.time(20); - if (RUN_OCL_IMPL) { ocl::oclMat oclSrc(src), oclDst(srcSize, type); @@ -217,9 +278,8 @@ PERF_TEST_P(SobelFixture, Sobel, typedef Size_MatType ScharrFixture; -PERF_TEST_P(ScharrFixture, Scharr, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(ScharrFixture, Scharr, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -230,12 +290,6 @@ PERF_TEST_P(ScharrFixture, Scharr, Mat src(srcSize, type), dst(srcSize, type); declare.in(src, WARMUP_RNG).out(dst); - if ((srcSize == OCL_SIZE_2000 && type == CV_8UC4) || - (srcSize == OCL_SIZE_4000 && type == CV_8UC1)) - declare.time(5.5); - else if (srcSize == OCL_SIZE_4000 && type == CV_8UC4) - declare.time(21); - if (RUN_OCL_IMPL) { ocl::oclMat oclSrc(src), oclDst(srcSize, type); @@ -244,7 +298,7 @@ PERF_TEST_P(ScharrFixture, Scharr, oclDst.download(dst); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 3e-3); } else if (RUN_PLAIN_IMPL) { @@ -258,15 +312,14 @@ PERF_TEST_P(ScharrFixture, Scharr, ///////////// GaussianBlur //////////////////////// -typedef Size_MatType GaussianBlurFixture; +typedef FilterFixture GaussianBlurFixture; -PERF_TEST_P(GaussianBlurFixture, GaussianBlur, - ::testing::Combine(::testing::Values(OCL_SIZE_1000, OCL_SIZE_2000), - OCL_PERF_ENUM(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4))) +OCL_PERF_TEST_P(GaussianBlurFixture, GaussianBlur, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5, 7))) { - const Size_MatType_t params = GetParam(); + const FilterParams params = GetParam(); const Size srcSize = get<0>(params); - const int type = get<1>(params), ksize = 7; + const int type = get<1>(params), ksize = get<2>(params); checkDeviceMaxMemoryAllocSize(srcSize, type); @@ -297,15 +350,14 @@ PERF_TEST_P(GaussianBlurFixture, GaussianBlur, ///////////// filter2D//////////////////////// -typedef Size_MatType filter2DFixture; +typedef FilterFixture Filter2DFixture; -PERF_TEST_P(filter2DFixture, filter2D, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(Filter2DFixture, Filter2D, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5))) { - const Size_MatType_t params = GetParam(); + const FilterParams params = GetParam(); const Size srcSize = get<0>(params); - const int type = get<1>(params), ksize = 3; + const int type = get<1>(params), ksize = get<2>(params); checkDeviceMaxMemoryAllocSize(srcSize, type); @@ -313,9 +365,6 @@ PERF_TEST_P(filter2DFixture, filter2D, declare.in(src, WARMUP_RNG).in(kernel).out(dst); randu(kernel, -3.0, 3.0); - if (srcSize == OCL_SIZE_4000 && type == CV_8UC4) - declare.time(8); - if (RUN_OCL_IMPL) { ocl::oclMat oclSrc(src), oclDst(srcSize, type), oclKernel(kernel); @@ -324,13 +373,13 @@ PERF_TEST_P(filter2DFixture, filter2D, oclDst.download(dst); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 3e-2); } else if (RUN_PLAIN_IMPL) { TEST_CYCLE() cv::filter2D(src, dst, -1, kernel); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 1e-2); } else OCL_PERF_ELSE @@ -338,28 +387,22 @@ PERF_TEST_P(filter2DFixture, filter2D, ///////////// Bilateral//////////////////////// -typedef Size_MatType BilateralFixture; +typedef TestBaseWithParam BilateralFixture; -PERF_TEST_P(BilateralFixture, Bilateral, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC3))) +OCL_PERF_TEST_P(BilateralFixture, Bilateral, OCL_TEST_SIZES) { - const Size_MatType_t params = GetParam(); - const Size srcSize = get<0>(params); - const int type = get<1>(params), d = 7; + const Size srcSize = GetParam(); + const int d = 7; const double sigmacolor = 50.0, sigmaspace = 50.0; - checkDeviceMaxMemoryAllocSize(srcSize, type); + checkDeviceMaxMemoryAllocSize(srcSize, CV_8UC1); - Mat src(srcSize, type), dst(srcSize, type); + Mat src(srcSize, CV_8UC1), dst(srcSize, CV_8UC1); declare.in(src, WARMUP_RNG).out(dst); - if (srcSize == OCL_SIZE_4000) - declare.time(type == CV_8UC3 ? 8 : 4.5); - if (RUN_OCL_IMPL) { - ocl::oclMat oclSrc(src), oclDst(srcSize, type); + ocl::oclMat oclSrc(src), oclDst(srcSize, CV_8UC1); OCL_TEST_CYCLE() cv::ocl::bilateralFilter(oclSrc, oclDst, d, sigmacolor, sigmaspace); @@ -377,37 +420,35 @@ PERF_TEST_P(BilateralFixture, Bilateral, OCL_PERF_ELSE } -///////////// adaptiveBilateral//////////////////////// +///////////// MedianBlur//////////////////////// -typedef Size_MatType adaptiveBilateralFixture; +typedef tuple MedianBlurParams; +typedef TestBaseWithParam MedianBlurFixture; -PERF_TEST_P(adaptiveBilateralFixture, DISABLED_adaptiveBilateral, - ::testing::Combine(::testing::Values(OCL_SIZE_1000), OCL_PERF_ENUM(CV_8UC1, CV_8UC3))) +OCL_PERF_TEST_P(MedianBlurFixture, Bilateral, ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(3, 5))) { - const Size_MatType_t params = GetParam(); + MedianBlurParams params = GetParam(); const Size srcSize = get<0>(params); - const int type = get<1>(params); - const double sigmaspace = 10.0; - Size ksize(9, 9); + const int ksize = get<1>(params); - checkDeviceMaxMemoryAllocSize(srcSize, type); + checkDeviceMaxMemoryAllocSize(srcSize, CV_8UC1); - Mat src(srcSize, type), dst(srcSize, type); + Mat src(srcSize, CV_8UC1), dst(srcSize, CV_8UC1); declare.in(src, WARMUP_RNG).out(dst); if (RUN_OCL_IMPL) { - ocl::oclMat oclSrc(src), oclDst(srcSize, type); + ocl::oclMat oclSrc(src), oclDst(srcSize, CV_8UC1); - OCL_TEST_CYCLE() cv::ocl::adaptiveBilateralFilter(oclSrc, oclDst, ksize, sigmaspace); + OCL_TEST_CYCLE() cv::ocl::medianFilter(oclSrc, oclDst, ksize); oclDst.download(dst); - SANITY_CHECK(dst, 1.0); + SANITY_CHECK(dst); } else if (RUN_PLAIN_IMPL) { - TEST_CYCLE() cv::adaptiveBilateralFilter(src, dst, ksize, sigmaspace); + TEST_CYCLE() cv::medianBlur(src, dst, ksize); SANITY_CHECK(dst); } diff --git a/modules/ocl/perf/perf_gemm.cpp b/modules/ocl/perf/perf_gemm.cpp index 4dcd5d4d6..32492ad97 100644 --- a/modules/ocl/perf/perf_gemm.cpp +++ b/modules/ocl/perf/perf_gemm.cpp @@ -46,16 +46,21 @@ #include "perf_precomp.hpp" using namespace perf; +using std::tr1::get; ///////////// gemm //////////////////////// -typedef TestBaseWithParam gemmFixture; +typedef Size_MatType GemmFixture; #ifdef HAVE_CLAMDBLAS -PERF_TEST_P(gemmFixture, gemm, ::testing::Values(OCL_SIZE_1000, OCL_SIZE_2000)) +OCL_PERF_TEST_P(GemmFixture, Gemm, ::testing::Combine( + ::testing::Values(Size(1000, 1000), Size(1500, 1500)), + ::testing::Values((int)cv::GEMM_1_T, (int)cv::GEMM_1_T | (int)cv::GEMM_2_T))) { - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); Mat src1(srcSize, CV_32FC1), src2(srcSize, CV_32FC1), src3(srcSize, CV_32FC1), dst(srcSize, CV_32FC1); @@ -69,7 +74,7 @@ PERF_TEST_P(gemmFixture, gemm, ::testing::Values(OCL_SIZE_1000, OCL_SIZE_2000)) ocl::oclMat oclSrc1(src1), oclSrc2(src2), oclSrc3(src3), oclDst(srcSize, CV_32FC1); - OCL_TEST_CYCLE() cv::ocl::gemm(oclSrc1, oclSrc2, 1.0, oclSrc3, 1.0, oclDst); + OCL_TEST_CYCLE() cv::ocl::gemm(oclSrc1, oclSrc2, 1.0, oclSrc3, 1.0, oclDst, type); oclDst.download(dst); @@ -77,7 +82,7 @@ PERF_TEST_P(gemmFixture, gemm, ::testing::Values(OCL_SIZE_1000, OCL_SIZE_2000)) } else if (RUN_PLAIN_IMPL) { - TEST_CYCLE() cv::gemm(src1, src2, 1.0, src3, 1.0, dst); + TEST_CYCLE() cv::gemm(src1, src2, 1.0, src3, 1.0, dst, type); SANITY_CHECK(dst, 0.01); } diff --git a/modules/ocl/perf/perf_gftt.cpp b/modules/ocl/perf/perf_gftt.cpp index af24c3489..688d24316 100644 --- a/modules/ocl/perf/perf_gftt.cpp +++ b/modules/ocl/perf/perf_gftt.cpp @@ -52,22 +52,21 @@ using std::tr1::get; ///////////// GoodFeaturesToTrack //////////////////////// -typedef tuple GoodFeaturesToTrackParams; +typedef tuple GoodFeaturesToTrackParams; typedef TestBaseWithParam GoodFeaturesToTrackFixture; -PERF_TEST_P(GoodFeaturesToTrackFixture, GoodFeaturesToTrack, - ::testing::Combine(::testing::Values(string("gpu/opticalflow/rubberwhale1.png"), - string("gpu/stereobm/aloe-L.png")), - ::testing::Range(0.0, 4.0, 3.0))) +OCL_PERF_TEST_P(GoodFeaturesToTrackFixture, GoodFeaturesToTrack, + ::testing::Combine(OCL_PERF_ENUM(String("gpu/opticalflow/rubberwhale1.png")), + OCL_PERF_ENUM(0.0, 3.0), testing::Bool())) { + const GoodFeaturesToTrackParams params = GetParam(); + const String fileName = get<0>(params); + const double minDistance = get<1>(params), qualityLevel = 0.01; + const bool harrisDetector = get<2>(params); + const int maxCorners = 1000; - const GoodFeaturesToTrackParams param = GetParam(); - const string fileName = getDataPath(get<0>(param)); - const int maxCorners = 2000; - const double qualityLevel = 0.01, minDistance = get<1>(param); - - Mat frame = imread(fileName, IMREAD_GRAYSCALE); - ASSERT_TRUE(!frame.empty()) << "no input image"; + Mat frame = imread(getDataPath(fileName), IMREAD_GRAYSCALE); + ASSERT_FALSE(frame.empty()) << "no input image"; vector pts_gold; declare.in(frame); @@ -75,7 +74,8 @@ PERF_TEST_P(GoodFeaturesToTrackFixture, GoodFeaturesToTrack, if (RUN_OCL_IMPL) { ocl::oclMat oclFrame(frame), pts_oclmat; - ocl::GoodFeaturesToTrackDetector_OCL detector(maxCorners, qualityLevel, minDistance); + ocl::GoodFeaturesToTrackDetector_OCL detector(maxCorners, qualityLevel, minDistance, 3, + harrisDetector); OCL_TEST_CYCLE() detector(oclFrame, pts_oclmat); @@ -86,7 +86,7 @@ PERF_TEST_P(GoodFeaturesToTrackFixture, GoodFeaturesToTrack, else if (RUN_PLAIN_IMPL) { TEST_CYCLE() cv::goodFeaturesToTrack(frame, pts_gold, - maxCorners, qualityLevel, minDistance); + maxCorners, qualityLevel, minDistance, noArray(), 3, harrisDetector); SANITY_CHECK(pts_gold); } diff --git a/modules/ocl/perf/perf_haar.cpp b/modules/ocl/perf/perf_haar.cpp index 1a97e908c..e70641d06 100644 --- a/modules/ocl/perf/perf_haar.cpp +++ b/modules/ocl/perf/perf_haar.cpp @@ -46,8 +46,13 @@ #include "perf_precomp.hpp" using namespace perf; +using namespace std; +using namespace cv; +using std::tr1::make_tuple; +using std::tr1::get; ///////////// Haar //////////////////////// + PERF_TEST(HaarFixture, Haar) { vector faces; @@ -84,23 +89,16 @@ PERF_TEST(HaarFixture, Haar) OCL_PERF_ELSE } -using namespace std; -using namespace cv; -using namespace perf; -using std::tr1::make_tuple; -using std::tr1::get; +typedef std::tr1::tuple Cascade_Image_MinSize_t; +typedef perf::TestBaseWithParam Cascade_Image_MinSize; -typedef std::tr1::tuple OCL_Cascade_Image_MinSize_t; -typedef perf::TestBaseWithParam OCL_Cascade_Image_MinSize; - -PERF_TEST_P( OCL_Cascade_Image_MinSize, CascadeClassifier, - testing::Combine( - testing::Values( string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt.xml"), - string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt2.xml") ), - testing::Values( string("cv/shared/lena.png"), - string("cv/cascadeandhog/images/bttf301.png")/*, - string("cv/cascadeandhog/images/class57.png")*/ ), - testing::Values(30, 64, 90) ) ) +OCL_PERF_TEST_P(Cascade_Image_MinSize, DISABLED_CascadeClassifier, + testing::Combine(testing::Values( string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt.xml"), + string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt2.xml") ), + testing::Values(string("cv/shared/lena.png"), + string("cv/cascadeandhog/images/bttf301.png")/*, + string("cv/cascadeandhog/images/class57.png")*/ ), + testing::Values(30, 64, 90))) { const string cascasePath = get<0>(GetParam()); const string imagePath = get<1>(GetParam()); @@ -109,7 +107,7 @@ PERF_TEST_P( OCL_Cascade_Image_MinSize, CascadeClassifier, vector faces; Mat img = imread(getDataPath(imagePath), IMREAD_GRAYSCALE); - ASSERT_TRUE(!img.empty()) << "Can't load source image: " << getDataPath(imagePath); + ASSERT_FALSE(img.empty()) << "Can't load source image: " << getDataPath(imagePath); equalizeHist(img, img); declare.in(img); @@ -146,7 +144,7 @@ PERF_TEST_P( OCL_Cascade_Image_MinSize, CascadeClassifier, else OCL_PERF_ELSE - //sort(faces.begin(), faces.end(), comparators::RectLess()); - SANITY_CHECK_NOTHING();//(faces, min_size/5); - // using SANITY_CHECK_NOTHING() since OCL and PLAIN version may find different faces number + //sort(faces.begin(), faces.end(), comparators::RectLess()); + SANITY_CHECK_NOTHING();//(faces, min_size/5); + // using SANITY_CHECK_NOTHING() since OCL and PLAIN version may find different faces number } diff --git a/modules/ocl/perf/perf_hog.cpp b/modules/ocl/perf/perf_hog.cpp index 2a6731117..a65769a26 100644 --- a/modules/ocl/perf/perf_hog.cpp +++ b/modules/ocl/perf/perf_hog.cpp @@ -43,7 +43,9 @@ // the use of this software, even if advised of the possibility of such damage. // //M*/ + #include "perf_precomp.hpp" +#include using namespace perf; @@ -66,7 +68,7 @@ struct RectLess : } }; -PERF_TEST(HOGFixture, HOG) +OCL_PERF_TEST(HOGFixture, HOG) { Mat src = imread(getDataPath("gpu/hog/road.png"), cv::IMREAD_GRAYSCALE); ASSERT_TRUE(!src.empty()) << "can't open input image road.png"; diff --git a/modules/ocl/perf/perf_imgproc.cpp b/modules/ocl/perf/perf_imgproc.cpp index e0ffe56e7..85f2c5528 100644 --- a/modules/ocl/perf/perf_imgproc.cpp +++ b/modules/ocl/perf/perf_imgproc.cpp @@ -51,9 +51,9 @@ using std::tr1::get; ///////////// equalizeHist //////////////////////// -typedef TestBaseWithParam equalizeHistFixture; +typedef TestBaseWithParam EqualizeHistFixture; -PERF_TEST_P(equalizeHistFixture, equalizeHist, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(EqualizeHistFixture, EqualizeHist, OCL_TEST_SIZES) { const Size srcSize = GetParam(); const double eps = 1 + DBL_EPSILON; @@ -83,16 +83,13 @@ PERF_TEST_P(equalizeHistFixture, equalizeHist, OCL_TYPICAL_MAT_SIZES) /////////// CopyMakeBorder ////////////////////// -CV_ENUM(Border, BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT, - BORDER_WRAP, BORDER_REFLECT_101) +CV_ENUM(Border, BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101) typedef tuple CopyMakeBorderParamType; typedef TestBaseWithParam CopyMakeBorderFixture; -PERF_TEST_P(CopyMakeBorderFixture, CopyMakeBorder, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4), - Border::all())) +OCL_PERF_TEST_P(CopyMakeBorderFixture, CopyMakeBorder, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, Border::all())) { const CopyMakeBorderParamType params = GetParam(); const Size srcSize = get<0>(params); @@ -125,11 +122,10 @@ PERF_TEST_P(CopyMakeBorderFixture, CopyMakeBorder, ///////////// cornerMinEigenVal //////////////////////// -typedef Size_MatType cornerMinEigenValFixture; +typedef Size_MatType CornerMinEigenValFixture; -PERF_TEST_P(cornerMinEigenValFixture, cornerMinEigenVal, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +OCL_PERF_TEST_P(CornerMinEigenValFixture, CornerMinEigenVal, + ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -165,11 +161,10 @@ PERF_TEST_P(cornerMinEigenValFixture, cornerMinEigenVal, ///////////// cornerHarris //////////////////////// -typedef Size_MatType cornerHarrisFixture; +typedef Size_MatType CornerHarrisFixture; -PERF_TEST_P(cornerHarrisFixture, cornerHarris, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +OCL_PERF_TEST_P(CornerHarrisFixture, CornerHarris, + ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -202,11 +197,14 @@ PERF_TEST_P(cornerHarrisFixture, cornerHarris, ///////////// integral //////////////////////// -typedef TestBaseWithParam integralFixture; +typedef tuple IntegralParams; +typedef TestBaseWithParam IntegralFixture; -PERF_TEST_P(integralFixture, integral, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(IntegralFixture, Integral1, ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32S, CV_32F))) { - const Size srcSize = GetParam(); + const IntegralParams params = GetParam(); + const Size srcSize = get<0>(params); + const int sdepth = get<1>(params); Mat src(srcSize, CV_8UC1), dst; declare.in(src, WARMUP_RNG); @@ -215,17 +213,17 @@ PERF_TEST_P(integralFixture, integral, OCL_TYPICAL_MAT_SIZES) { ocl::oclMat oclSrc(src), oclDst; - OCL_TEST_CYCLE() cv::ocl::integral(oclSrc, oclDst); + OCL_TEST_CYCLE() cv::ocl::integral(oclSrc, oclDst, sdepth); oclDst.download(dst); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE); } else if (RUN_PLAIN_IMPL) { - TEST_CYCLE() cv::integral(src, dst); + TEST_CYCLE() cv::integral(src, dst, sdepth); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE); } else OCL_PERF_ELSE @@ -233,15 +231,13 @@ PERF_TEST_P(integralFixture, integral, OCL_TYPICAL_MAT_SIZES) ///////////// threshold//////////////////////// -CV_ENUM(ThreshType, THRESH_BINARY, THRESH_TOZERO_INV) +CV_ENUM(ThreshType, THRESH_BINARY, THRESH_BINARY_INV, THRESH_TRUNC, THRESH_TOZERO_INV) typedef tuple ThreshParams; typedef TestBaseWithParam ThreshFixture; -PERF_TEST_P(ThreshFixture, threshold, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4, CV_16SC1, CV_16SC4, CV_32FC1), - ThreshType::all())) +OCL_PERF_TEST_P(ThreshFixture, Threshold, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, ThreshType::all())) { const ThreshParams params = GetParam(); const Size srcSize = get<0>(params); @@ -463,9 +459,9 @@ static void meanShiftFiltering_(const Mat &src_roi, Mat &dst_roi, int sp, int sr } } -typedef TestBaseWithParam meanShiftFilteringFixture; +typedef TestBaseWithParam MeanShiftFilteringFixture; -PERF_TEST_P(meanShiftFilteringFixture, meanShiftFiltering, +PERF_TEST_P(MeanShiftFilteringFixture, MeanShiftFiltering, OCL_TYPICAL_MAT_SIZES) { const Size srcSize = GetParam(); @@ -556,9 +552,9 @@ static void meanShiftProc_(const Mat &src_roi, Mat &dst_roi, Mat &dstCoor_roi, i } -typedef TestBaseWithParam meanShiftProcFixture; +typedef TestBaseWithParam MeanShiftProcFixture; -PERF_TEST_P(meanShiftProcFixture, meanShiftProc, +PERF_TEST_P(MeanShiftProcFixture, MeanShiftProc, OCL_TYPICAL_MAT_SIZES) { const Size srcSize = GetParam(); @@ -598,7 +594,7 @@ PERF_TEST_P(meanShiftProcFixture, meanShiftProc, typedef TestBaseWithParam CLAHEFixture; -PERF_TEST_P(CLAHEFixture, CLAHE, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(CLAHEFixture, CLAHE, OCL_TEST_SIZES) { const Size srcSize = GetParam(); const string impl = getSelectedImpl(); @@ -632,9 +628,9 @@ PERF_TEST_P(CLAHEFixture, CLAHE, OCL_TYPICAL_MAT_SIZES) OCL_PERF_ELSE } -///////////// columnSum//////////////////////// +///////////// ColumnSum//////////////////////// -typedef TestBaseWithParam columnSumFixture; +typedef TestBaseWithParam ColumnSumFixture; static void columnSumPerfTest(const Mat & src, Mat & dst) { @@ -646,7 +642,7 @@ static void columnSumPerfTest(const Mat & src, Mat & dst) dst.at(i, j) = dst.at(i - 1 , j) + src.at(i , j); } -PERF_TEST_P(columnSumFixture, columnSum, OCL_TYPICAL_MAT_SIZES) +PERF_TEST_P(ColumnSumFixture, ColumnSum, OCL_TYPICAL_MAT_SIZES) { const Size srcSize = GetParam(); @@ -680,8 +676,8 @@ PERF_TEST_P(columnSumFixture, columnSum, OCL_TYPICAL_MAT_SIZES) CV_ENUM(DistType, NORM_L1, NORM_L2SQR) -typedef tuple distanceToCentersParameters; -typedef TestBaseWithParam distanceToCentersFixture; +typedef tuple DistanceToCentersParams; +typedef TestBaseWithParam DistanceToCentersFixture; static void distanceToCentersPerfTest(Mat& src, Mat& centers, Mat& dists, Mat& labels, int distType) { @@ -706,10 +702,11 @@ static void distanceToCentersPerfTest(Mat& src, Mat& centers, Mat& dists, Mat& l Mat(labels_v).copyTo(labels); } -PERF_TEST_P(distanceToCentersFixture, distanceToCenters, ::testing::Combine(::testing::Values(cv::Size(256,256), cv::Size(512,512)), DistType::all()) ) +PERF_TEST_P(DistanceToCentersFixture, DistanceToCenters, ::testing::Combine(::testing::Values(cv::Size(256,256), cv::Size(512,512)), DistType::all()) ) { - Size size = get<0>(GetParam()); - int distType = get<1>(GetParam()); + const DistanceToCentersParams params = GetParam(); + Size size = get<0>(params); + int distType = get<1>(params); Mat src(size, CV_32FC1), centers(size, CV_32FC1); Mat dists(src.rows, 1, CV_32FC1), labels(src.rows, 1, CV_32SC1); diff --git a/modules/ocl/perf/perf_imgwarp.cpp b/modules/ocl/perf/perf_imgwarp.cpp index e768d6621..fe863fa70 100644 --- a/modules/ocl/perf/perf_imgwarp.cpp +++ b/modules/ocl/perf/perf_imgwarp.cpp @@ -51,11 +51,13 @@ using std::tr1::get; ///////////// WarpAffine //////////////////////// -typedef Size_MatType WarpAffineFixture; +CV_ENUM(InterType, INTER_NEAREST, INTER_LINEAR) -PERF_TEST_P(WarpAffineFixture, WarpAffine, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +typedef tuple WarpAffineParams; +typedef TestBaseWithParam WarpAffineFixture; + +OCL_PERF_TEST_P(WarpAffineFixture, WarpAffine, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134, InterType::all())) { static const double coeffs[2][3] = { @@ -63,11 +65,12 @@ PERF_TEST_P(WarpAffineFixture, WarpAffine, { sin(CV_PI / 6), cos(CV_PI / 6), -100.0 } }; Mat M(2, 3, CV_64F, (void *)coeffs); - const int interpolation = INTER_NEAREST; - const Size_MatType_t params = GetParam(); + const WarpAffineParams params = GetParam(); const Size srcSize = get<0>(params); - const int type = get<1>(params); + const int type = get<1>(params), interpolation = get<2>(params); + + checkDeviceMaxMemoryAllocSize(srcSize, type); Mat src(srcSize, type), dst(srcSize, type); declare.in(src, WARMUP_RNG).out(dst); @@ -80,13 +83,13 @@ PERF_TEST_P(WarpAffineFixture, WarpAffine, oclDst.download(dst); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 5e-4); } else if (RUN_PLAIN_IMPL) { TEST_CYCLE() cv::warpAffine(src, dst, M, srcSize, interpolation); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 5e-4); } else OCL_PERF_ELSE @@ -94,11 +97,11 @@ PERF_TEST_P(WarpAffineFixture, WarpAffine, ///////////// WarpPerspective //////////////////////// -typedef Size_MatType WarpPerspectiveFixture; +typedef WarpAffineParams WarpPerspectiveParams; +typedef TestBaseWithParam WarpPerspectiveFixture; -PERF_TEST_P(WarpPerspectiveFixture, WarpPerspective, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(WarpPerspectiveFixture, WarpPerspective, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134, InterType::all())) { static const double coeffs[3][3] = { @@ -107,15 +110,15 @@ PERF_TEST_P(WarpPerspectiveFixture, WarpPerspective, {0.0, 0.0, 1.0} }; Mat M(3, 3, CV_64F, (void *)coeffs); - const int interpolation = INTER_LINEAR; - const Size_MatType_t params = GetParam(); + const WarpPerspectiveParams params = GetParam(); const Size srcSize = get<0>(params); - const int type = get<1>(params); + const int type = get<1>(params), interpolation = get<2>(params); + + checkDeviceMaxMemoryAllocSize(srcSize, type); Mat src(srcSize, type), dst(srcSize, type); - declare.in(src, WARMUP_RNG).out(dst) - .time(srcSize == OCL_SIZE_4000 ? 18 : srcSize == OCL_SIZE_2000 ? 5 : 2); + declare.in(src, WARMUP_RNG).out(dst); if (RUN_OCL_IMPL) { @@ -125,32 +128,28 @@ PERF_TEST_P(WarpPerspectiveFixture, WarpPerspective, oclDst.download(dst); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 5e-4); } else if (RUN_PLAIN_IMPL) { TEST_CYCLE() cv::warpPerspective(src, dst, M, srcSize, interpolation); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 5e-4); } else OCL_PERF_ELSE } -///////////// resize //////////////////////// +///////////// Resize //////////////////////// -CV_ENUM(resizeInterType, INTER_NEAREST, INTER_LINEAR) +typedef tuple ResizeParams; +typedef TestBaseWithParam ResizeFixture; -typedef tuple resizeParams; -typedef TestBaseWithParam resizeFixture; - -PERF_TEST_P(resizeFixture, resize, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4), - resizeInterType::all(), - ::testing::Values(0.5, 2.0))) +OCL_PERF_TEST_P(ResizeFixture, Resize, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134, + InterType::all(), ::testing::Values(0.5, 2.0))) { - const resizeParams params = GetParam(); + const ResizeParams params = GetParam(); const Size srcSize = get<0>(params); const int type = get<1>(params), interType = get<2>(params); double scale = get<3>(params); @@ -162,8 +161,6 @@ PERF_TEST_P(resizeFixture, resize, Mat src(srcSize, type), dst; dst.create(dstSize, type); declare.in(src, WARMUP_RNG).out(dst); - if (interType == INTER_LINEAR && type == CV_8UC4 && OCL_SIZE_4000 == srcSize) - declare.time(11); if (RUN_OCL_IMPL) { @@ -185,15 +182,13 @@ PERF_TEST_P(resizeFixture, resize, OCL_PERF_ELSE } -typedef tuple resizeAreaParams; -typedef TestBaseWithParam resizeAreaFixture; +typedef tuple ResizeAreaParams; +typedef TestBaseWithParam ResizeAreaFixture; -PERF_TEST_P(resizeAreaFixture, resize, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4, CV_32FC1, CV_32FC4), - ::testing::Values(0.3, 0.5, 0.6))) +OCL_PERF_TEST_P(ResizeAreaFixture, Resize, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES_134, ::testing::Values(0.3, 0.5, 0.6))) { - const resizeAreaParams params = GetParam(); + const ResizeAreaParams params = GetParam(); const Size srcSize = get<0>(params); const int type = get<1>(params); double scale = get<2>(params); @@ -225,19 +220,15 @@ PERF_TEST_P(resizeAreaFixture, resize, OCL_PERF_ELSE } -///////////// remap//////////////////////// +///////////// Remap //////////////////////// -CV_ENUM(RemapInterType, INTER_NEAREST, INTER_LINEAR) +typedef tuple RemapParams; +typedef TestBaseWithParam RemapFixture; -typedef tuple remapParams; -typedef TestBaseWithParam remapFixture; - -PERF_TEST_P(remapFixture, remap, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4), - RemapInterType::all())) +OCL_PERF_TEST_P(RemapFixture, Remap, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, InterType::all())) { - const remapParams params = GetParam(); + const RemapParams params = GetParam(); const Size srcSize = get<0>(params); const int type = get<1>(params), interpolation = get<2>(params); @@ -287,7 +278,7 @@ PERF_TEST_P(remapFixture, remap, } -///////////// buildWarpPerspectiveMaps //////////////////////// +///////////// BuildWarpPerspectiveMaps //////////////////////// static void buildWarpPerspectiveMaps(const Mat &M, bool inverse, Size dsize, Mat &xmap, Mat &ymap) { @@ -323,9 +314,9 @@ static void buildWarpPerspectiveMaps(const Mat &M, bool inverse, Size dsize, Mat } } -typedef TestBaseWithParam buildWarpPerspectiveMapsFixture; +typedef TestBaseWithParam BuildWarpPerspectiveMapsFixture; -PERF_TEST_P(buildWarpPerspectiveMapsFixture, Inverse, OCL_TYPICAL_MAT_SIZES) +PERF_TEST_P(BuildWarpPerspectiveMapsFixture, Inverse, OCL_TYPICAL_MAT_SIZES) { static const double coeffs[3][3] = { diff --git a/modules/ocl/perf/perf_match_template.cpp b/modules/ocl/perf/perf_match_template.cpp index 561a978ae..236ae1750 100644 --- a/modules/ocl/perf/perf_match_template.cpp +++ b/modules/ocl/perf/perf_match_template.cpp @@ -43,6 +43,7 @@ // the use of this software, even if advised of the possibility of such damage. // //M*/ + #include "perf_precomp.hpp" using namespace perf; @@ -53,8 +54,8 @@ using std::tr1::get; typedef Size_MatType CV_TM_CCORRFixture; -PERF_TEST_P(CV_TM_CCORRFixture, matchTemplate, - ::testing::Combine(::testing::Values(OCL_SIZE_1000, OCL_SIZE_2000), +OCL_PERF_TEST_P(CV_TM_CCORRFixture, matchTemplate, + ::testing::Combine(::testing::Values(Size(1000, 1000), Size(2000, 2000)), OCL_PERF_ENUM(CV_32FC1, CV_32FC4))) { const Size_MatType_t params = GetParam(); @@ -66,7 +67,7 @@ PERF_TEST_P(CV_TM_CCORRFixture, matchTemplate, Mat dst(dstSize, CV_32F); randu(src, 0.0f, 1.0f); randu(templ, 0.0f, 1.0f); - declare.time(srcSize == OCL_SIZE_2000 ? 20 : 6).in(src, templ).out(dst); + declare.in(src, templ).out(dst); if (RUN_OCL_IMPL) { @@ -90,7 +91,8 @@ PERF_TEST_P(CV_TM_CCORRFixture, matchTemplate, typedef TestBaseWithParam CV_TM_CCORR_NORMEDFixture; -PERF_TEST_P(CV_TM_CCORR_NORMEDFixture, matchTemplate, OCL_TYPICAL_MAT_SIZES) +OCL_PERF_TEST_P(CV_TM_CCORR_NORMEDFixture, matchTemplate, + ::testing::Values(Size(1000, 1000), Size(2000, 2000), Size(4000, 4000))) { const Size srcSize = GetParam(), templSize(5, 5); @@ -125,7 +127,7 @@ CV_ENUM(MethodType, TM_SQDIFF, TM_SQDIFF_NORMED, TM_CCORR, TM_CCORR_NORMED, TM_C typedef std::tr1::tuple ImgSize_TmplSize_Method_t; typedef TestBaseWithParam ImgSize_TmplSize_Method; -PERF_TEST_P(ImgSize_TmplSize_Method, MatchTemplate, +OCL_PERF_TEST_P(ImgSize_TmplSize_Method, MatchTemplate, ::testing::Combine( testing::Values(szSmall128, cv::Size(320, 240), cv::Size(640, 480), cv::Size(800, 600), diff --git a/modules/ocl/perf/perf_matrix_operation.cpp b/modules/ocl/perf/perf_matrix_operation.cpp index 5ca322e22..434f134ca 100644 --- a/modules/ocl/perf/perf_matrix_operation.cpp +++ b/modules/ocl/perf/perf_matrix_operation.cpp @@ -53,9 +53,7 @@ using std::tr1::get; typedef Size_MatType ConvertToFixture; -PERF_TEST_P(ConvertToFixture, ConvertTo, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(ConvertToFixture, ConvertTo, ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -92,11 +90,9 @@ PERF_TEST_P(ConvertToFixture, ConvertTo, ///////////// copyTo//////////////////////// -typedef Size_MatType copyToFixture; +typedef Size_MatType CopyToFixture; -PERF_TEST_P(copyToFixture, copyTo, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(CopyToFixture, CopyTo, ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -127,11 +123,9 @@ PERF_TEST_P(copyToFixture, copyTo, ///////////// setTo//////////////////////// -typedef Size_MatType setToFixture; +typedef Size_MatType SetToFixture; -PERF_TEST_P(setToFixture, setTo, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(SetToFixture, SetTo, ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -164,16 +158,16 @@ PERF_TEST_P(setToFixture, setTo, /////////////////// upload /////////////////////////// -typedef tuple uploadParams; -typedef TestBaseWithParam uploadFixture; +typedef tuple UploadParams; +typedef TestBaseWithParam UploadFixture; -PERF_TEST_P(uploadFixture, upload, +PERF_TEST_P(UploadFixture, Upload, testing::Combine( OCL_TYPICAL_MAT_SIZES, testing::Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F), testing::Range(1, 5))) { - const uploadParams params = GetParam(); + const UploadParams params = GetParam(); const Size srcSize = get<0>(params); const int depth = get<1>(params), cn = get<2>(params); const int type = CV_MAKE_TYPE(depth, cn); @@ -201,15 +195,15 @@ PERF_TEST_P(uploadFixture, upload, /////////////////// download /////////////////////////// -typedef TestBaseWithParam downloadFixture; +typedef TestBaseWithParam DownloadFixture; -PERF_TEST_P(downloadFixture, download, +PERF_TEST_P(DownloadFixture, Download, testing::Combine( OCL_TYPICAL_MAT_SIZES, testing::Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F), testing::Range(1, 5))) { - const uploadParams params = GetParam(); + const UploadParams params = GetParam(); const Size srcSize = get<0>(params); const int depth = get<1>(params), cn = get<2>(params); const int type = CV_MAKE_TYPE(depth, cn); diff --git a/modules/ocl/perf/perf_moments.cpp b/modules/ocl/perf/perf_moments.cpp index 631031ecb..d6a9b5c20 100644 --- a/modules/ocl/perf/perf_moments.cpp +++ b/modules/ocl/perf/perf_moments.cpp @@ -55,22 +55,19 @@ using namespace cvtest; using namespace testing; using namespace std; - ///////////// Moments //////////////////////// -//*! performance of image -typedef tuple MomentsParamType; -typedef TestBaseWithParam MomentsFixture; -PERF_TEST_P(MomentsFixture, Moments, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_16SC1, CV_16UC1, CV_32FC1), ::testing::Bool())) +typedef tuple MomentsParams; +typedef TestBaseWithParam MomentsFixture; + +OCL_PERF_TEST_P(MomentsFixture, Moments, + ::testing::Combine(OCL_TEST_SIZES, ::testing::Bool())) { - const MomentsParamType params = GetParam(); + const MomentsParams params = GetParam(); const Size srcSize = get<0>(params); - const int type = get<1>(params); - const bool binaryImage = get<2>(params); + const bool binaryImage = get<1>(params); - Mat src(srcSize, type), dst(7, 1, CV_64F); + Mat src(srcSize, CV_8UC1), dst(7, 1, CV_64F); randu(src, 0, 255); cv::Moments mom; @@ -85,6 +82,7 @@ PERF_TEST_P(MomentsFixture, Moments, } else OCL_PERF_ELSE + cv::HuMoments(mom, dst); SANITY_CHECK(dst, 2e-1); } diff --git a/modules/ocl/perf/perf_norm.cpp b/modules/ocl/perf/perf_norm.cpp index ff49eb4ed..abeb93cf1 100644 --- a/modules/ocl/perf/perf_norm.cpp +++ b/modules/ocl/perf/perf_norm.cpp @@ -51,18 +51,21 @@ using std::tr1::get; ///////////// norm//////////////////////// -typedef tuple normParams; -typedef TestBaseWithParam normFixture; +CV_ENUM(NormType, NORM_INF, NORM_L1, NORM_L2) -PERF_TEST_P(normFixture, norm, testing::Combine( - OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +typedef std::tr1::tuple NormParams; +typedef TestBaseWithParam NormFixture; + +OCL_PERF_TEST_P(NormFixture, Norm, + ::testing::Combine(OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3), + OCL_TEST_TYPES, NormType::all())) { - const normParams params = GetParam(); + const NormParams params = GetParam(); const Size srcSize = get<0>(params); const int type = get<1>(params); - double value = 0.0; - const double eps = CV_MAT_DEPTH(type) == CV_8U ? DBL_EPSILON : 1e-3; + const int normType = get<2>(params); + perf::ERROR_TYPE errorType = type != NORM_INF ? ERROR_RELATIVE : ERROR_ABSOLUTE; + double eps = 1e-5, value; Mat src1(srcSize, type), src2(srcSize, type); declare.in(src1, src2, WARMUP_RNG); @@ -71,15 +74,15 @@ PERF_TEST_P(normFixture, norm, testing::Combine( { ocl::oclMat oclSrc1(src1), oclSrc2(src2); - OCL_TEST_CYCLE() value = cv::ocl::norm(oclSrc1, oclSrc2, NORM_INF); + OCL_TEST_CYCLE() value = cv::ocl::norm(oclSrc1, oclSrc2, normType); - SANITY_CHECK(value, eps); + SANITY_CHECK(value, eps, errorType); } else if (RUN_PLAIN_IMPL) { - TEST_CYCLE() value = cv::norm(src1, src2, NORM_INF); + TEST_CYCLE() value = cv::norm(src1, src2, normType); - SANITY_CHECK(value); + SANITY_CHECK(value, eps, errorType); } else OCL_PERF_ELSE diff --git a/modules/ocl/perf/perf_opticalflow.cpp b/modules/ocl/perf/perf_opticalflow.cpp index bc1761b49..3711e87d8 100644 --- a/modules/ocl/perf/perf_opticalflow.cpp +++ b/modules/ocl/perf/perf_opticalflow.cpp @@ -52,55 +52,28 @@ using std::tr1::get; using std::tr1::tuple; using std::tr1::make_tuple; -CV_ENUM(LoadMode, IMREAD_GRAYSCALE, IMREAD_COLOR) +typedef tuple PyrLKOpticalFlowParamType; +typedef TestBaseWithParam PyrLKOpticalFlowFixture; -typedef tuple > PyrLKOpticalFlowParamType; -typedef TestBaseWithParam PyrLKOpticalFlowFixture; - -PERF_TEST_P(PyrLKOpticalFlowFixture, - PyrLKOpticalFlow, - ::testing::Combine( - ::testing::Values(1000, 2000, 4000), - ::testing::Values( - make_tuple - ( - string("gpu/opticalflow/rubberwhale1.png"), - string("gpu/opticalflow/rubberwhale2.png"), - LoadMode(IMREAD_COLOR) - ), - make_tuple - ( - string("gpu/stereobm/aloe-L.png"), - string("gpu/stereobm/aloe-R.png"), - LoadMode(IMREAD_GRAYSCALE) - ) - ) - ) - ) +OCL_PERF_TEST_P(PyrLKOpticalFlowFixture, + PyrLKOpticalFlow, ::testing::Values(1000, 2000, 4000)) { - PyrLKOpticalFlowParamType params = GetParam(); - tuple fileParam = get<1>(params); - const int pointsCount = get<0>(params); - const int openMode = static_cast(get<2>(fileParam)); - const string fileName0 = get<0>(fileParam), fileName1 = get<1>(fileParam); - Mat frame0 = imread(getDataPath(fileName0), openMode); - Mat frame1 = imread(getDataPath(fileName1), openMode); + const int pointsCount = GetParam(); + + const string fileName0 = "gpu/opticalflow/rubberwhale1.png", + fileName1 = "gpu/opticalflow/rubberwhale2.png"; + Mat frame0 = imread(getDataPath(fileName0), cv::IMREAD_GRAYSCALE); + Mat frame1 = imread(getDataPath(fileName1), cv::IMREAD_GRAYSCALE); declare.in(frame0, frame1); ASSERT_FALSE(frame0.empty()) << "can't load " << fileName0; ASSERT_FALSE(frame1.empty()) << "can't load " << fileName1; - Mat grayFrame; - if (openMode == IMREAD_COLOR) - cvtColor(frame0, grayFrame, COLOR_BGR2GRAY); - else - grayFrame = frame0; - vector pts, nextPts; vector status; vector err; - goodFeaturesToTrack(grayFrame, pts, pointsCount, 0.01, 0.0); + goodFeaturesToTrack(frame0, pts, pointsCount, 0.01, 0.0); Mat ptsMat(1, static_cast(pts.size()), CV_32FC2, (void *)&pts[0]); if (RUN_PLAIN_IMPL) @@ -178,12 +151,11 @@ CV_ENUM(farneFlagType, 0, OPTFLOW_FARNEBACK_GAUSSIAN) typedef tuple, farneFlagType, bool> FarnebackOpticalFlowParams; typedef TestBaseWithParam FarnebackOpticalFlowFixture; -PERF_TEST_P(FarnebackOpticalFlowFixture, FarnebackOpticalFlow, - ::testing::Combine( - ::testing::Values(make_tuple(5, 1.1), - make_tuple(7, 1.5)), - farneFlagType::all(), - ::testing::Bool())) +OCL_PERF_TEST_P(FarnebackOpticalFlowFixture, FarnebackOpticalFlow, + ::testing::Combine( + ::testing::Values(make_tuple(5, 1.1), + make_tuple(7, 1.5)), + farneFlagType::all(), ::testing::Bool())) { Mat frame0 = imread(getDataPath("gpu/opticalflow/rubberwhale1.png"), cv::IMREAD_GRAYSCALE); ASSERT_FALSE(frame0.empty()) << "can't load rubberwhale1.png"; diff --git a/modules/ocl/perf/perf_precomp.hpp b/modules/ocl/perf/perf_precomp.hpp index a8663df99..5f92be820 100644 --- a/modules/ocl/perf/perf_precomp.hpp +++ b/modules/ocl/perf/perf_precomp.hpp @@ -70,8 +70,7 @@ #include "opencv2/ocl/ocl.hpp" #include "opencv2/ts/ts.hpp" -using namespace std; -using namespace cv; +// TODO remove it #define OCL_SIZE_1000 Size(1000, 1000) #define OCL_SIZE_2000 Size(2000, 2000) @@ -79,6 +78,19 @@ using namespace cv; #define OCL_TYPICAL_MAT_SIZES ::testing::Values(OCL_SIZE_1000, OCL_SIZE_2000, OCL_SIZE_4000) +using namespace std; +using namespace cv; + +#define OCL_SIZE_1 szVGA +#define OCL_SIZE_2 sz720p +#define OCL_SIZE_3 sz1080p +#define OCL_SIZE_4 sz2160p + +#define OCL_TEST_SIZES ::testing::Values(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3, OCL_SIZE_4) +#define OCL_TEST_TYPES ::testing::Values(CV_8UC1, CV_32FC1, CV_8UC4, CV_32FC4) +#define OCL_TEST_TYPES_14 OCL_TEST_TYPES +#define OCL_TEST_TYPES_134 ::testing::Values(CV_8UC1, CV_32FC1, CV_8UC3, CV_32FC3, CV_8UC4, CV_32FC4) + #define OCL_PERF_ENUM(type, ...) ::testing::Values(type, ## __VA_ARGS__ ) #define IMPL_OCL "ocl" @@ -103,6 +115,31 @@ using namespace cv; CV_TEST_FAIL_NO_IMPL(); #endif +#define OCL_PERF_TEST(fixture, name) \ + class OCL##_##fixture##_##name : \ + public ::perf::TestBase \ + { \ + public: \ + OCL##_##fixture##_##name() { } \ + protected: \ + virtual void PerfTestBody(); \ + }; \ + TEST_F(OCL##_##fixture##_##name, name) { RunPerfTestBody(); } \ + void OCL##_##fixture##_##name::PerfTestBody() + +#define OCL_PERF_TEST_P(fixture, name, params) \ + class OCL##_##fixture##_##name : \ + public fixture \ + { \ + public: \ + OCL##_##fixture##_##name() { } \ + protected: \ + virtual void PerfTestBody(); \ + }; \ + TEST_P(OCL##_##fixture##_##name, name) { RunPerfTestBody(); } \ + INSTANTIATE_TEST_CASE_P(/*none*/, OCL##_##fixture##_##name, params); \ + void OCL##_##fixture##_##name::PerfTestBody() + #define OCL_TEST_CYCLE_N(n) for(declare.iterations(n); startTimer(), next(); cv::ocl::finish(), stopTimer()) #define OCL_TEST_CYCLE() for(; startTimer(), next(); cv::ocl::finish(), stopTimer()) #define OCL_TEST_CYCLE_MULTIRUN(runsNum) for(declare.runs(runsNum); startTimer(), next(); stopTimer()) for(int r = 0; r < runsNum; cv::ocl::finish(), ++r) diff --git a/modules/ocl/perf/perf_pyramid.cpp b/modules/ocl/perf/perf_pyramid.cpp index 820dd6062..b4ca25356 100644 --- a/modules/ocl/perf/perf_pyramid.cpp +++ b/modules/ocl/perf/perf_pyramid.cpp @@ -51,11 +51,10 @@ using std::tr1::get; ///////////// pyrDown ////////////////////// -typedef Size_MatType pyrDownFixture; +typedef Size_MatType PyrDownFixture; -PERF_TEST_P(pyrDownFixture, pyrDown, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(PyrDownFixture, PyrDown, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -77,7 +76,7 @@ PERF_TEST_P(pyrDownFixture, pyrDown, oclDst.download(dst); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 5e-4); } else if (RUN_PLAIN_IMPL) { @@ -91,11 +90,10 @@ PERF_TEST_P(pyrDownFixture, pyrDown, ///////////// pyrUp //////////////////////// -typedef Size_MatType pyrUpFixture; +typedef Size_MatType PyrUpFixture; -PERF_TEST_P(pyrUpFixture, pyrUp, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_8UC4))) +OCL_PERF_TEST_P(PyrUpFixture, PyrUp, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) { const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); @@ -117,7 +115,7 @@ PERF_TEST_P(pyrUpFixture, pyrUp, oclDst.download(dst); - SANITY_CHECK(dst); + SANITY_CHECK(dst, 5e-4); } else if (RUN_PLAIN_IMPL) { diff --git a/modules/ocl/perf/perf_split_merge.cpp b/modules/ocl/perf/perf_split_merge.cpp index ecfc49e33..ab3f82f58 100644 --- a/modules/ocl/perf/perf_split_merge.cpp +++ b/modules/ocl/perf/perf_split_merge.cpp @@ -51,21 +51,22 @@ using std::tr1::get; ///////////// Merge//////////////////////// -typedef Size_MatType MergeFixture; +typedef tuple MergeParams; +typedef TestBaseWithParam MergeFixture; -PERF_TEST_P(MergeFixture, Merge, - ::testing::Combine(::testing::Values(OCL_SIZE_1000, OCL_SIZE_2000), - OCL_PERF_ENUM(CV_8U, CV_32F))) +OCL_PERF_TEST_P(MergeFixture, Merge, + ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8U, CV_32F), + OCL_PERF_ENUM(2, 3))) { - const Size_MatType_t params = GetParam(); + const MergeParams params = GetParam(); const Size srcSize = get<0>(params); - const int depth = get<1>(params), channels = 3; - const int dstType = CV_MAKE_TYPE(depth, channels); + const int depth = get<1>(params), cn = get<2>(params), + dtype = CV_MAKE_TYPE(depth, cn); - checkDeviceMaxMemoryAllocSize(srcSize, dstType); + checkDeviceMaxMemoryAllocSize(srcSize, dtype); - Mat dst(srcSize, dstType); - vector src(channels); + Mat dst(srcSize, dtype); + vector src(cn); for (vector::iterator i = src.begin(), end = src.end(); i != end; ++i) { i->create(srcSize, CV_MAKE_TYPE(depth, 1)); @@ -75,7 +76,7 @@ PERF_TEST_P(MergeFixture, Merge, if (RUN_OCL_IMPL) { - ocl::oclMat oclDst(srcSize, dstType); + ocl::oclMat oclDst(srcSize, dtype); vector oclSrc(src.size()); for (vector::size_type i = 0, end = src.size(); i < end; ++i) oclSrc[i] = src[i]; @@ -98,49 +99,69 @@ PERF_TEST_P(MergeFixture, Merge, ///////////// Split//////////////////////// -typedef Size_MatType SplitFixture; +typedef MergeParams SplitParams; +typedef TestBaseWithParam SplitFixture; -PERF_TEST_P(SplitFixture, Split, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8U, CV_32F))) +OCL_PERF_TEST_P(SplitFixture, Split, + ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8U, CV_32F), + OCL_PERF_ENUM(2, 3))) { - const Size_MatType_t params = GetParam(); + const SplitParams params = GetParam(); const Size srcSize = get<0>(params); - const int depth = get<1>(params), channels = 3; - const int type = CV_MAKE_TYPE(depth, channels); + const int depth = get<1>(params), cn = get<2>(params); + const int type = CV_MAKE_TYPE(depth, cn); checkDeviceMaxMemoryAllocSize(srcSize, type); Mat src(srcSize, type); + Mat dst0, dst1, dst2; declare.in(src, WARMUP_RNG); + ASSERT_TRUE(cn == 3 || cn == 2); + if (RUN_OCL_IMPL) { ocl::oclMat oclSrc(src); - vector oclDst(channels, ocl::oclMat(srcSize, CV_MAKE_TYPE(depth, 1))); + vector oclDst(cn); + oclDst[0] = ocl::oclMat(srcSize, depth); + oclDst[1] = ocl::oclMat(srcSize, depth); + if (cn == 3) + oclDst[2] = ocl::oclMat(srcSize, depth); OCL_TEST_CYCLE() cv::ocl::split(oclSrc, oclDst); - ASSERT_EQ(3, channels); - Mat dst0, dst1, dst2; oclDst[0].download(dst0); oclDst[1].download(dst1); - oclDst[2].download(dst2); - SANITY_CHECK(dst0); - SANITY_CHECK(dst1); - SANITY_CHECK(dst2); + if (cn == 3) + oclDst[2].download(dst2); } else if (RUN_PLAIN_IMPL) { - vector dst(channels, Mat(srcSize, CV_MAKE_TYPE(depth, 1))); + vector dst(cn); + dst[0] = Mat(srcSize, depth); + dst[1] = Mat(srcSize, depth); + if (cn == 3) + dst[2] = Mat(srcSize, depth); + TEST_CYCLE() cv::split(src, dst); - ASSERT_EQ(3, channels); - Mat & dst0 = dst[0], & dst1 = dst[1], & dst2 = dst[2]; + dst0 = dst[0]; + dst1 = dst[1]; + if (cn == 3) + dst2 = dst[2]; + } + else + OCL_PERF_ELSE + + if (cn == 2) + { + SANITY_CHECK(dst0); + SANITY_CHECK(dst1); + } + else if (cn == 3) + { SANITY_CHECK(dst0); SANITY_CHECK(dst1); SANITY_CHECK(dst2); } - else - OCL_PERF_ELSE } From fbc69e444b67b1643f62be065ca7b1e777abfed6 Mon Sep 17 00:00:00 2001 From: Cody Rigney Date: Mon, 24 Feb 2014 14:15:21 -0500 Subject: [PATCH 02/10] Added NEON optimizations for LK optical flow (Intrinsics). --- modules/video/src/lkpyramid.cpp | 252 ++++++++++++++++++++++++++++++++ modules/video/src/precomp.hpp | 1 + 2 files changed, 253 insertions(+) diff --git a/modules/video/src/lkpyramid.cpp b/modules/video/src/lkpyramid.cpp index 291cb86a2..e807db2d1 100644 --- a/modules/video/src/lkpyramid.cpp +++ b/modules/video/src/lkpyramid.cpp @@ -87,6 +87,30 @@ static void calcSharrDeriv(const cv::Mat& src, cv::Mat& dst) _mm_store_si128((__m128i*)(trow1 + x), t1); } #endif + +#if CV_NEON + + const uint16x8_t q8 = vdupq_n_u16(3); + const uint8x8_t d18 = vdup_n_u8(10); + + const int16x8_t q8i = vdupq_n_s16(3); + const int16x8_t q9 = vdupq_n_s16(10); + + for( ; x <= colsn - 8; x += 8) + { + uint8x8_t d0 = vld1_u8((const uint8_t*)&srow0[x]); + uint8x8_t d1 = vld1_u8((const uint8_t*)&srow1[x]); + uint8x8_t d2 = vld1_u8((const uint8_t*)&srow2[x]); + uint16x8_t q4 = vaddl_u8(d0, d2); + uint16x8_t q11 = vsubl_u8(d2, d0); + uint16x8_t q5 = vmulq_u16(q4, q8); + uint16x8_t q6 = vmull_u8(d1, d18); + uint16x8_t q10 = vaddq_u16(q6, q5); + vst1q_u16((uint16_t*)&trow0[x], q10); + vst1q_u16((uint16_t*)&trow1[x], q11); + + } +#endif for( ; x < colsn; x++ ) { int t0 = (srow0[x] + srow2[x])*3 + srow1[x]*10; @@ -123,6 +147,35 @@ static void calcSharrDeriv(const cv::Mat& src, cv::Mat& dst) _mm_storeu_si128((__m128i*)(drow + x*2 + 8), t0); } #endif + +#if CV_NEON + for( ; x <= colsn - 8; x += 8 ) + { + + int16x8_t q0 = vld1q_s16((const int16_t*)&trow0[x+cn]); + int16x8_t q1 = vld1q_s16((const int16_t*)&trow0[x-cn]); + int16x8_t q2 = vld1q_s16((const int16_t*)&trow1[x+cn]); + int16x8_t q3 = vld1q_s16((const int16_t*)&trow1[x-cn]); + int16x8_t q5 = vsubq_s16(q0, q1); + int16x8_t q6 = vaddq_s16(q2, q3); + int16x8_t q4 = vld1q_s16((const int16_t*)&trow1[x]); + int16x8_t q7 = vmulq_s16(q6, q8i); + int16x8_t q10 = vmulq_s16(q4, q9); + int16x8_t q11 = vaddq_s16(q7, q10); + int16x4_t d22 = vget_low_s16(q11); + int16x4_t d23 = vget_high_s16(q11); + int16x4_t d11 = vget_high_s16(q5); + int16x4_t d10 = vget_low_s16(q5); + int16x4x2_t q5x2, q11x2; + q5x2.val[0] = d10; q5x2.val[1] = d22; + q11x2.val[0] = d11; q11x2.val[1] = d23; + vst2_s16((int16_t*)&drow[x*2], q5x2); + vst2_s16((int16_t*)&drow[(x*2)+8], q11x2); + + } + +#endif + for( ; x < colsn; x++ ) { deriv_type t0 = (deriv_type)(trow0[x+cn] - trow0[x-cn]); @@ -226,6 +279,21 @@ void cv::detail::LKTrackerInvoker::operator()(const Range& range) const __m128 qA11 = _mm_setzero_ps(), qA12 = _mm_setzero_ps(), qA22 = _mm_setzero_ps(); #endif +#if CV_NEON + + int CV_DECL_ALIGNED(16) nA11[] = {0, 0, 0, 0}, nA12[] = {0, 0, 0, 0}, nA22[] = {0, 0, 0, 0}; + const int shifter1 = -(W_BITS - 5); //negative so it shifts right + const int shifter2 = -(W_BITS); + + const int16x4_t d26 = vdup_n_s16((int16_t)iw00); + const int16x4_t d27 = vdup_n_s16((int16_t)iw01); + const int16x4_t d28 = vdup_n_s16((int16_t)iw10); + const int16x4_t d29 = vdup_n_s16((int16_t)iw11); + const int32x4_t q11 = vdupq_n_s32((int32_t)shifter1); + const int32x4_t q12 = vdupq_n_s32((int32_t)shifter2); + +#endif + // extract the patch from the first image, compute covariation matrix of derivatives int x, y; for( y = 0; y < winSize.height; y++ ) @@ -279,6 +347,92 @@ void cv::detail::LKTrackerInvoker::operator()(const Range& range) const } #endif +#if CV_NEON + + for( ; x <= winSize.width*cn - 4; x += 4, dsrc += 4*2, dIptr += 4*2 ) + { + + uint8x8_t d0 = vld1_u8(&src[x]); + uint8x8_t d2 = vld1_u8(&src[x+cn]); + uint16x8_t q0 = vmovl_u8(d0); + uint16x8_t q1 = vmovl_u8(d2); + + int32x4_t q5 = vmull_s16(vget_low_s16(vreinterpretq_s16_u16(q0)), d26); + int32x4_t q6 = vmull_s16(vget_low_s16(vreinterpretq_s16_u16(q1)), d27); + + uint8x8_t d4 = vld1_u8(&src[x + stepI]); + uint8x8_t d6 = vld1_u8(&src[x + stepI + cn]); + uint16x8_t q2 = vmovl_u8(d4); + uint16x8_t q3 = vmovl_u8(d6); + + int32x4_t q7 = vmull_s16(vget_low_s16(vreinterpretq_s16_u16(q2)), d28); + int32x4_t q8 = vmull_s16(vget_low_s16(vreinterpretq_s16_u16(q3)), d29); + + q5 = vaddq_s32(q5, q6); + q7 = vaddq_s32(q7, q8); + q5 = vaddq_s32(q5, q7); + + int16x4x2_t d0d1 = vld2_s16(dsrc); + int16x4x2_t d2d3 = vld2_s16(&dsrc[cn2]); + + q5 = vqrshlq_s32(q5, q11); + + int32x4_t q4 = vmull_s16(d0d1.val[0], d26); + q6 = vmull_s16(d0d1.val[1], d26); + + int16x4_t nd0 = vmovn_s32(q5); + + q7 = vmull_s16(d2d3.val[0], d27); + q8 = vmull_s16(d2d3.val[1], d27); + + vst1_s16(&Iptr[x], nd0); + + int16x4x2_t d4d5 = vld2_s16(&dsrc[dstep]); + int16x4x2_t d6d7 = vld2_s16(&dsrc[dstep+cn2]); + + q4 = vaddq_s32(q4, q7); + q6 = vaddq_s32(q6, q8); + + q7 = vmull_s16(d4d5.val[0], d28); + int32x4_t nq0 = vmull_s16(d4d5.val[1], d28); + q8 = vmull_s16(d6d7.val[0], d29); + int32x4_t q15 = vmull_s16(d6d7.val[1], d29); + + q7 = vaddq_s32(q7, q8); + nq0 = vaddq_s32(nq0, q15); + + q4 = vaddq_s32(q4, q7); + q6 = vaddq_s32(q6, nq0); + + int32x4_t nq1 = vld1q_s32(nA12); + int32x4_t nq2 = vld1q_s32(nA22); + nq0 = vld1q_s32(nA11); + + q4 = vqrshlq_s32(q4, q12); + q6 = vqrshlq_s32(q6, q12); + + q7 = vmulq_s32(q4, q4); + q8 = vmulq_s32(q4, q6); + q15 = vmulq_s32(q6, q6); + + nq0 = vaddq_s32(nq0, q7); + nq1 = vaddq_s32(nq1, q8); + nq2 = vaddq_s32(nq2, q15); + + vst1q_s32(nA11, nq0); + vst1q_s32(nA12, nq1); + vst1q_s32(nA22, nq2); + + int16x4_t d8 = vmovn_s32(q4); + int16x4_t d12 = vmovn_s32(q6); + + int16x4x2_t d8d12; + d8d12.val[0] = d8; d8d12.val[1] = d12; + vst2_s16(dIptr, d8d12); + } + +#endif + for( ; x < winSize.width*cn; x++, dsrc += 2, dIptr += 2 ) { int ival = CV_DESCALE(src[x]*iw00 + src[x+cn]*iw01 + @@ -308,6 +462,12 @@ void cv::detail::LKTrackerInvoker::operator()(const Range& range) const A22 += A22buf[0] + A22buf[1] + A22buf[2] + A22buf[3]; #endif +#if CV_NEON + A11 += (float)(nA11[0] + nA11[1] + nA11[2] + nA11[3]); + A12 += (float)(nA12[0] + nA12[1] + nA12[2] + nA12[3]); + A22 += (float)(nA22[0] + nA22[1] + nA22[2] + nA22[3]); +#endif + A11 *= FLT_SCALE; A12 *= FLT_SCALE; A22 *= FLT_SCALE; @@ -357,6 +517,17 @@ void cv::detail::LKTrackerInvoker::operator()(const Range& range) const __m128 qb0 = _mm_setzero_ps(), qb1 = _mm_setzero_ps(); #endif +#if CV_NEON + + int CV_DECL_ALIGNED(16) nB1[] = {0,0,0,0}, nB2[] = {0,0,0,0}; + + const int16x4_t d26_2 = vdup_n_s16((int16_t)iw00); + const int16x4_t d27_2 = vdup_n_s16((int16_t)iw01); + const int16x4_t d28_2 = vdup_n_s16((int16_t)iw10); + const int16x4_t d29_2 = vdup_n_s16((int16_t)iw11); + +#endif + for( y = 0; y < winSize.height; y++ ) { const uchar* Jptr = (const uchar*)J.data + (y + inextPt.y)*stepJ + inextPt.x*cn; @@ -400,6 +571,80 @@ void cv::detail::LKTrackerInvoker::operator()(const Range& range) const } #endif +#if CV_NEON + + for( ; x <= winSize.width*cn - 8; x += 8, dIptr += 8*2 ) + { + + uint8x8_t d0 = vld1_u8(&Jptr[x]); + uint8x8_t d2 = vld1_u8(&Jptr[x+cn]); + uint8x8_t d4 = vld1_u8(&Jptr[x+stepJ]); + uint8x8_t d6 = vld1_u8(&Jptr[x+stepJ+cn]); + + uint16x8_t q0 = vmovl_u8(d0); + uint16x8_t q1 = vmovl_u8(d2); + uint16x8_t q2 = vmovl_u8(d4); + uint16x8_t q3 = vmovl_u8(d6); + + int32x4_t nq4 = vmull_s16(vget_low_s16(vreinterpretq_s16_u16(q0)), d26_2); + int32x4_t nq5 = vmull_s16(vget_high_s16(vreinterpretq_s16_u16(q0)), d26_2); + + int32x4_t nq6 = vmull_s16(vget_low_s16(vreinterpretq_s16_u16(q1)), d27_2); + int32x4_t nq7 = vmull_s16(vget_high_s16(vreinterpretq_s16_u16(q1)), d27_2); + + int32x4_t nq8 = vmull_s16(vget_low_s16(vreinterpretq_s16_u16(q2)), d28_2); + int32x4_t nq9 = vmull_s16(vget_high_s16(vreinterpretq_s16_u16(q2)), d28_2); + + int32x4_t nq10 = vmull_s16(vget_low_s16(vreinterpretq_s16_u16(q3)), d29_2); + int32x4_t nq11 = vmull_s16(vget_high_s16(vreinterpretq_s16_u16(q3)), d29_2); + + nq4 = vaddq_s32(nq4, nq6); + nq5 = vaddq_s32(nq5, nq7); + nq8 = vaddq_s32(nq8, nq10); + nq9 = vaddq_s32(nq9, nq11); + + int16x8_t q6 = vld1q_s16(&Iptr[x]); + + nq4 = vaddq_s32(nq4, nq8); + nq5 = vaddq_s32(nq5, nq9); + + nq8 = vmovl_s16(vget_high_s16(q6)); + nq6 = vmovl_s16(vget_low_s16(q6)); + + nq4 = vqrshlq_s32(nq4, q11); + nq5 = vqrshlq_s32(nq5, q11); + + int16x8x2_t q0q1 = vld2q_s16(dIptr); + nq11 = vld1q_s32(nB1); + int32x4_t nq15 = vld1q_s32(nB2); + + nq4 = vsubq_s32(nq4, nq6); + nq5 = vsubq_s32(nq5, nq8); + + int32x4_t nq2 = vmovl_s16(vget_low_s16(q0q1.val[0])); + int32x4_t nq3 = vmovl_s16(vget_high_s16(q0q1.val[0])); + + nq7 = vmovl_s16(vget_low_s16(q0q1.val[1])); + nq8 = vmovl_s16(vget_high_s16(q0q1.val[1])); + + nq9 = vmulq_s32(nq4, nq2); + nq10 = vmulq_s32(nq5, nq3); + + nq4 = vmulq_s32(nq4, nq7); + nq5 = vmulq_s32(nq5, nq8); + + nq9 = vaddq_s32(nq9, nq10); + nq4 = vaddq_s32(nq4, nq5); + + nq11 = vaddq_s32(nq11, nq9); + nq15 = vaddq_s32(nq15, nq4); + + vst1q_s32(nB1, nq11); + vst1q_s32(nB2, nq15); + } + +#endif + for( ; x < winSize.width*cn; x++, dIptr += 2 ) { int diff = CV_DESCALE(Jptr[x]*iw00 + Jptr[x+cn]*iw01 + @@ -417,6 +662,13 @@ void cv::detail::LKTrackerInvoker::operator()(const Range& range) const b2 += bbuf[1] + bbuf[3]; #endif +#if CV_NEON + + b1 += (float)(nB1[0] + nB1[1] + nB1[2] + nB1[3]); + b2 += (float)(nB2[0] + nB2[1] + nB2[2] + nB2[3]); + +#endif + b1 *= FLT_SCALE; b2 *= FLT_SCALE; diff --git a/modules/video/src/precomp.hpp b/modules/video/src/precomp.hpp index 9b6077f1e..ced9a3761 100644 --- a/modules/video/src/precomp.hpp +++ b/modules/video/src/precomp.hpp @@ -49,6 +49,7 @@ #include "opencv2/video/background_segm.hpp" #include "opencv2/imgproc/imgproc_c.h" #include "opencv2/core/internal.hpp" +#include "opencv2/core/core.hpp" #include From 767b28f2e3bf75883edaf79e5a4de5566a1926e2 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sat, 1 Mar 2014 13:13:24 +0400 Subject: [PATCH 03/10] tests --- modules/ocl/perf/perf_arithm.cpp | 194 +++--------- modules/ocl/perf/perf_brute_force_matcher.cpp | 33 ++- modules/ocl/perf/perf_fft.cpp | 3 - modules/ocl/perf/perf_gemm.cpp | 12 +- modules/ocl/perf/perf_hog.cpp | 2 +- modules/ocl/perf/perf_imgproc.cpp | 20 +- modules/ocl/perf/perf_imgwarp.cpp | 3 - modules/ocl/perf/perf_kalman.cpp | 4 +- modules/ocl/perf/perf_match_template.cpp | 3 +- modules/ocl/perf/perf_ml.cpp | 14 +- modules/ocl/perf/perf_norm.cpp | 89 ------ modules/ocl/perf/perf_opticalflow.cpp | 7 +- modules/ocl/perf/perf_stat.cpp | 276 ++++++++++++++++++ 13 files changed, 364 insertions(+), 296 deletions(-) delete mode 100644 modules/ocl/perf/perf_norm.cpp create mode 100644 modules/ocl/perf/perf_stat.cpp diff --git a/modules/ocl/perf/perf_arithm.cpp b/modules/ocl/perf/perf_arithm.cpp index 7c194ae16..3faedfcc7 100644 --- a/modules/ocl/perf/perf_arithm.cpp +++ b/modules/ocl/perf/perf_arithm.cpp @@ -60,15 +60,14 @@ OCL_PERF_TEST_P(LUTFixture, LUT, // getting params const Size_MatType_t params = GetParam(); const Size srcSize = get<0>(params); - const int type = get<1>(params); + const int type = get<1>(params), cn = CV_MAT_CN(type); // creating src data - Mat src(srcSize, CV_8UC1), lut(1, 256, type); + Mat src(srcSize, CV_8UC(cn)), lut(1, 256, type); int dstType = CV_MAKETYPE(lut.depth(), src.channels()); Mat dst(srcSize, dstType); - randu(lut, 0, 2); - declare.in(src, WARMUP_RNG).in(lut).out(dst); + declare.in(src, lut, WARMUP_RNG).out(dst); // select implementation if (RUN_OCL_IMPL) @@ -564,158 +563,6 @@ OCL_PERF_TEST_P(FlipFixture, Flip, OCL_PERF_ELSE } -///////////// MinMax //////////////////////// - -typedef Size_MatType MinMaxFixture; - -PERF_TEST_P(MinMaxFixture, MinMax, - ::testing::Combine(OCL_TYPICAL_MAT_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) -{ - const Size_MatType_t params = GetParam(); - const Size srcSize = get<0>(params); - const int type = get<1>(params); - - Mat src(srcSize, type); - declare.in(src, WARMUP_RNG); - - double min_val = std::numeric_limits::max(), - max_val = std::numeric_limits::min(); - - if (RUN_OCL_IMPL) - { - ocl::oclMat oclSrc(src); - - OCL_TEST_CYCLE() cv::ocl::minMax(oclSrc, &min_val, &max_val); - - ASSERT_GE(max_val, min_val); - SANITY_CHECK(min_val); - SANITY_CHECK(max_val); - } - else if (RUN_PLAIN_IMPL) - { - Point min_loc, max_loc; - - TEST_CYCLE() cv::minMaxLoc(src, &min_val, &max_val, &min_loc, &max_loc); - - ASSERT_GE(max_val, min_val); - SANITY_CHECK(min_val); - SANITY_CHECK(max_val); - } - else - OCL_PERF_ELSE -} - -///////////// MinMaxLoc //////////////////////// - -typedef Size_MatType MinMaxLocFixture; - -OCL_PERF_TEST_P(MinMaxLocFixture, MinMaxLoc, - ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) -{ - const Size_MatType_t params = GetParam(); - const Size srcSize = get<0>(params); - const int type = get<1>(params); - - Mat src(srcSize, type); - randu(src, 0, 1); - declare.in(src); - - double min_val = 0.0, max_val = 0.0; - Point min_loc, max_loc; - - if (RUN_OCL_IMPL) - { - ocl::oclMat oclSrc(src); - - OCL_TEST_CYCLE() cv::ocl::minMaxLoc(oclSrc, &min_val, &max_val, &min_loc, &max_loc); - - ASSERT_GE(max_val, min_val); - SANITY_CHECK(min_val); - SANITY_CHECK(max_val); - } - else if (RUN_PLAIN_IMPL) - { - TEST_CYCLE() cv::minMaxLoc(src, &min_val, &max_val, &min_loc, &max_loc); - - ASSERT_GE(max_val, min_val); - SANITY_CHECK(min_val); - SANITY_CHECK(max_val); - } - else - OCL_PERF_ELSE -} - -///////////// Sum //////////////////////// - -typedef Size_MatType SumFixture; - -OCL_PERF_TEST_P(SumFixture, Sum, - ::testing::Combine(OCL_TEST_SIZES, - OCL_TEST_TYPES)) -{ - const Size_MatType_t params = GetParam(); - const Size srcSize = get<0>(params); - const int type = get<1>(params); - - Mat src(srcSize, type); - Scalar result; - randu(src, 0, 60); - declare.in(src); - - if (RUN_OCL_IMPL) - { - ocl::oclMat oclSrc(src); - - OCL_TEST_CYCLE() result = cv::ocl::sum(oclSrc); - - SANITY_CHECK(result, 1e-6, ERROR_RELATIVE); - } - else if (RUN_PLAIN_IMPL) - { - TEST_CYCLE() result = cv::sum(src); - - SANITY_CHECK(result, 1e-6, ERROR_RELATIVE); - } - else - OCL_PERF_ELSE -} - -///////////// countNonZero //////////////////////// - -typedef Size_MatType CountNonZeroFixture; - -OCL_PERF_TEST_P(CountNonZeroFixture, CountNonZero, - ::testing::Combine(OCL_TEST_SIZES, - OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) -{ - const Size_MatType_t params = GetParam(); - const Size srcSize = get<0>(params); - const int type = get<1>(params); - - Mat src(srcSize, type); - int result = 0; - randu(src, 0, 256); - declare.in(src); - - if (RUN_OCL_IMPL) - { - ocl::oclMat oclSrc(src); - - OCL_TEST_CYCLE() result = cv::ocl::countNonZero(oclSrc); - - SANITY_CHECK(result); - } - else if (RUN_PLAIN_IMPL) - { - TEST_CYCLE() result = cv::countNonZero(src); - - SANITY_CHECK(result); - } - else - OCL_PERF_ELSE -} - ///////////// Phase //////////////////////// typedef Size_MatType PhaseFixture; @@ -895,6 +742,41 @@ OCL_PERF_TEST_P(BitwiseNotFixture, Bitwise_not, OCL_PERF_ELSE } +///////////// SetIdentity //////////////////////// + +typedef Size_MatType SetIdentityFixture; + +OCL_PERF_TEST_P(SetIdentityFixture, SetIdentity, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) +{ + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); + + Mat src(srcSize, type); + Scalar s = Scalar::all(17); + declare.in(src, WARMUP_RNG).out(src); + + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc(src); + + OCL_TEST_CYCLE() cv::ocl::setIdentity(oclSrc, s); + + oclSrc.download(src); + + SANITY_CHECK(src); + } + else if (RUN_PLAIN_IMPL) + { + TEST_CYCLE() cv::setIdentity(src, s); + + SANITY_CHECK(src); + } + else + OCL_PERF_ELSE +} + ///////////// compare//////////////////////// CV_ENUM(CmpCode, CMP_LT, CMP_LE, CMP_EQ, CMP_NE, CMP_GE, CMP_GT) diff --git a/modules/ocl/perf/perf_brute_force_matcher.cpp b/modules/ocl/perf/perf_brute_force_matcher.cpp index 54e829f4b..b8ad81361 100644 --- a/modules/ocl/perf/perf_brute_force_matcher.cpp +++ b/modules/ocl/perf/perf_brute_force_matcher.cpp @@ -46,17 +46,22 @@ #include "perf_precomp.hpp" using namespace perf; +using std::tr1::get; //////////////////// BruteForceMatch ///////////////// -typedef TestBaseWithParam BruteForceMatcherFixture; +typedef Size_MatType BruteForceMatcherFixture; -OCL_PERF_TEST_P(BruteForceMatcherFixture, Match, OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3)) +OCL_PERF_TEST_P(BruteForceMatcherFixture, Match, + ::testing::Combine(OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3), + OCL_PERF_ENUM(MatType(CV_32FC1)))) { - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); vector matches; - Mat query(srcSize, CV_32FC1), train(srcSize, CV_32FC1); + Mat query(srcSize, type), train(srcSize, type); declare.in(query, train); randu(query, 0.0f, 1.0f); randu(train, 0.0f, 1.0f); @@ -82,12 +87,16 @@ OCL_PERF_TEST_P(BruteForceMatcherFixture, Match, OCL_PERF_ENUM(OCL_SIZE_1, OCL_S OCL_PERF_ELSE } -OCL_PERF_TEST_P(BruteForceMatcherFixture, KnnMatch, OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3)) +OCL_PERF_TEST_P(BruteForceMatcherFixture, KnnMatch, + ::testing::Combine(OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3), + OCL_PERF_ENUM(MatType(CV_32FC1)))) { - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); vector > matches(2); - Mat query(srcSize, CV_32F), train(srcSize, CV_32F); + Mat query(srcSize, type), train(srcSize, type); randu(query, 0.0f, 1.0f); randu(train, 0.0f, 1.0f); @@ -121,13 +130,17 @@ OCL_PERF_TEST_P(BruteForceMatcherFixture, KnnMatch, OCL_PERF_ENUM(OCL_SIZE_1, OC OCL_PERF_ELSE } -OCL_PERF_TEST_P(BruteForceMatcherFixture, RadiusMatch, OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3)) +OCL_PERF_TEST_P(BruteForceMatcherFixture, RadiusMatch, + ::testing::Combine(OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3), + OCL_PERF_ENUM(MatType(CV_32FC1)))) { - const Size srcSize = GetParam(); + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); const float max_distance = 2.0f; vector > matches(2); - Mat query(srcSize, CV_32FC1), train(srcSize, CV_32FC1); + Mat query(srcSize, type), train(srcSize, type); declare.in(query, train); randu(query, 0.0f, 1.0f); diff --git a/modules/ocl/perf/perf_fft.cpp b/modules/ocl/perf/perf_fft.cpp index 29e042ccd..95701b7da 100644 --- a/modules/ocl/perf/perf_fft.cpp +++ b/modules/ocl/perf/perf_fft.cpp @@ -71,9 +71,6 @@ OCL_PERF_TEST_P(DftFixture, Dft, ::testing::Combine(testing::Values(OCL_SIZE_1, randu(src, 0.0f, 1.0f); declare.in(src); - if (srcSize == OCL_SIZE_4000) - declare.time(7.4); - if (RUN_OCL_IMPL) { ocl::oclMat oclSrc(src), oclDst; diff --git a/modules/ocl/perf/perf_gemm.cpp b/modules/ocl/perf/perf_gemm.cpp index 32492ad97..69db3a099 100644 --- a/modules/ocl/perf/perf_gemm.cpp +++ b/modules/ocl/perf/perf_gemm.cpp @@ -47,28 +47,32 @@ using namespace perf; using std::tr1::get; +using std::tr1::tuple; ///////////// gemm //////////////////////// -typedef Size_MatType GemmFixture; - #ifdef HAVE_CLAMDBLAS +typedef tuple GemmParams; +typedef TestBaseWithParam GemmFixture; + OCL_PERF_TEST_P(GemmFixture, Gemm, ::testing::Combine( ::testing::Values(Size(1000, 1000), Size(1500, 1500)), ::testing::Values((int)cv::GEMM_1_T, (int)cv::GEMM_1_T | (int)cv::GEMM_2_T))) { - const Size_MatType_t params = GetParam(); + const GemmParams params = GetParam(); const Size srcSize = get<0>(params); const int type = get<1>(params); Mat src1(srcSize, CV_32FC1), src2(srcSize, CV_32FC1), src3(srcSize, CV_32FC1), dst(srcSize, CV_32FC1); - declare.in(src1, src2, src3).out(dst).time(srcSize == OCL_SIZE_2000 ? 65 : 8); + randu(src1, -10.0f, 10.0f); randu(src2, -10.0f, 10.0f); randu(src3, -10.0f, 10.0f); + declare.in(src1, src2, src3).out(dst); + if (RUN_OCL_IMPL) { ocl::oclMat oclSrc1(src1), oclSrc2(src2), diff --git a/modules/ocl/perf/perf_hog.cpp b/modules/ocl/perf/perf_hog.cpp index a65769a26..497b73d69 100644 --- a/modules/ocl/perf/perf_hog.cpp +++ b/modules/ocl/perf/perf_hog.cpp @@ -74,7 +74,7 @@ OCL_PERF_TEST(HOGFixture, HOG) ASSERT_TRUE(!src.empty()) << "can't open input image road.png"; vector found_locations; - declare.in(src).time(5); + declare.in(src); if (RUN_PLAIN_IMPL) { diff --git a/modules/ocl/perf/perf_imgproc.cpp b/modules/ocl/perf/perf_imgproc.cpp index 85f2c5528..e6f6a0582 100644 --- a/modules/ocl/perf/perf_imgproc.cpp +++ b/modules/ocl/perf/perf_imgproc.cpp @@ -133,8 +133,7 @@ OCL_PERF_TEST_P(CornerMinEigenValFixture, CornerMinEigenVal, const int blockSize = 7, apertureSize = 1 + 2 * 3; Mat src(srcSize, type), dst(srcSize, CV_32FC1); - declare.in(src, WARMUP_RNG).out(dst) - .time(srcSize == OCL_SIZE_4000 ? 20 : srcSize == OCL_SIZE_2000 ? 5 : 3); + declare.in(src, WARMUP_RNG).out(dst); const int depth = CV_MAT_DEPTH(type); const ERROR_TYPE errorType = depth == CV_8U ? ERROR_ABSOLUTE : ERROR_RELATIVE; @@ -172,8 +171,7 @@ OCL_PERF_TEST_P(CornerHarrisFixture, CornerHarris, Mat src(srcSize, type), dst(srcSize, CV_32FC1); randu(src, 0, 1); - declare.in(src).out(dst) - .time(srcSize == OCL_SIZE_4000 ? 20 : srcSize == OCL_SIZE_2000 ? 5 : 3); + declare.in(src).out(dst); if (RUN_OCL_IMPL) { @@ -469,9 +467,7 @@ PERF_TEST_P(MeanShiftFilteringFixture, MeanShiftFiltering, cv::TermCriteria crit(cv::TermCriteria::COUNT + cv::TermCriteria::EPS, 5, 1); Mat src(srcSize, CV_8UC4), dst(srcSize, CV_8UC4); - declare.in(src, WARMUP_RNG).out(dst) - .time(srcSize == OCL_SIZE_4000 ? - 56 : srcSize == OCL_SIZE_2000 ? 15 : 3.8); + declare.in(src, WARMUP_RNG).out(dst); if (RUN_PLAIN_IMPL) { @@ -562,9 +558,7 @@ PERF_TEST_P(MeanShiftProcFixture, MeanShiftProc, Mat src(srcSize, CV_8UC4), dst1(srcSize, CV_8UC4), dst2(srcSize, CV_16SC2); - declare.in(src, WARMUP_RNG).out(dst1, dst2) - .time(srcSize == OCL_SIZE_4000 ? - 56 : srcSize == OCL_SIZE_2000 ? 15 : 3.8);; + declare.in(src, WARMUP_RNG).out(dst1, dst2); if (RUN_PLAIN_IMPL) { @@ -603,9 +597,6 @@ OCL_PERF_TEST_P(CLAHEFixture, CLAHE, OCL_TEST_SIZES) const double clipLimit = 40.0; declare.in(src, WARMUP_RNG); - if (srcSize == OCL_SIZE_4000) - declare.time(11); - if (RUN_OCL_IMPL) { ocl::oclMat oclSrc(src), oclDst; @@ -649,9 +640,6 @@ PERF_TEST_P(ColumnSumFixture, ColumnSum, OCL_TYPICAL_MAT_SIZES) Mat src(srcSize, CV_32FC1), dst(srcSize, CV_32FC1); declare.in(src, WARMUP_RNG).out(dst); - if (srcSize == OCL_SIZE_4000) - declare.time(5); - if (RUN_OCL_IMPL) { ocl::oclMat oclSrc(src), oclDst(srcSize, CV_32FC1); diff --git a/modules/ocl/perf/perf_imgwarp.cpp b/modules/ocl/perf/perf_imgwarp.cpp index fe863fa70..56d55c03c 100644 --- a/modules/ocl/perf/perf_imgwarp.cpp +++ b/modules/ocl/perf/perf_imgwarp.cpp @@ -235,9 +235,6 @@ OCL_PERF_TEST_P(RemapFixture, Remap, Mat src(srcSize, type), dst(srcSize, type); declare.in(src, WARMUP_RNG).out(dst); - if (srcSize == OCL_SIZE_4000 && interpolation == INTER_LINEAR) - declare.time(9); - Mat xmap, ymap; xmap.create(srcSize, CV_32FC1); ymap.create(srcSize, CV_32FC1); diff --git a/modules/ocl/perf/perf_kalman.cpp b/modules/ocl/perf/perf_kalman.cpp index 946444ad9..e384fcd69 100644 --- a/modules/ocl/perf/perf_kalman.cpp +++ b/modules/ocl/perf/perf_kalman.cpp @@ -46,7 +46,7 @@ #include "perf_precomp.hpp" -#ifdef HAVE_CLAMDBLAS +//#ifdef HAVE_CLAMDBLAS using namespace perf; using namespace std; @@ -100,4 +100,4 @@ PERF_TEST_P(KalmanFilterFixture, KalmanFilter, SANITY_CHECK(statePre_); } -#endif // HAVE_CLAMDBLAS +//#endif // HAVE_CLAMDBLAS diff --git a/modules/ocl/perf/perf_match_template.cpp b/modules/ocl/perf/perf_match_template.cpp index 236ae1750..ae8c55719 100644 --- a/modules/ocl/perf/perf_match_template.cpp +++ b/modules/ocl/perf/perf_match_template.cpp @@ -99,8 +99,7 @@ OCL_PERF_TEST_P(CV_TM_CCORR_NORMEDFixture, matchTemplate, Mat src(srcSize, CV_8UC1), templ(templSize, CV_8UC1), dst; const Size dstSize(src.cols - templ.cols + 1, src.rows - templ.rows + 1); dst.create(dstSize, CV_8UC1); - declare.in(src, templ, WARMUP_RNG).out(dst) - .time(srcSize == OCL_SIZE_2000 ? 10 : srcSize == OCL_SIZE_4000 ? 23 : 2); + declare.in(src, templ, WARMUP_RNG).out(dst); if (RUN_OCL_IMPL) { diff --git a/modules/ocl/perf/perf_ml.cpp b/modules/ocl/perf/perf_ml.cpp index 13dcaa18f..680045673 100644 --- a/modules/ocl/perf/perf_ml.cpp +++ b/modules/ocl/perf/perf_ml.cpp @@ -55,7 +55,7 @@ static void genData(Mat& trainData, Size size, Mat& trainLabel = Mat().setTo(Sca trainData.create(size, CV_32FC1); randu(trainData, 1.0, 100.0); - if(nClasses != 0) + if (nClasses != 0) { trainLabel.create(size.height, 1, CV_8UC1); randu(trainLabel, 0, nClasses - 1); @@ -82,7 +82,7 @@ PERF_TEST_P(KNNFixture, KNN, genData(testData, size); Mat best_label; - if(RUN_PLAIN_IMPL) + if (RUN_PLAIN_IMPL) { TEST_CYCLE() { @@ -90,7 +90,8 @@ PERF_TEST_P(KNNFixture, KNN, knn_cpu.train(trainData, trainLabels); knn_cpu.find_nearest(testData, k, &best_label); } - }else if(RUN_OCL_IMPL) + } + else if (RUN_OCL_IMPL) { cv::ocl::oclMat best_label_ocl; cv::ocl::oclMat testdata; @@ -103,7 +104,8 @@ PERF_TEST_P(KNNFixture, KNN, knn_ocl.find_nearest(testdata, k, best_label_ocl); } best_label_ocl.download(best_label); - }else + } + else OCL_PERF_ELSE SANITY_CHECK(best_label); } @@ -188,7 +190,7 @@ PERF_TEST_P(SVMFixture, DISABLED_SVM, CvMat samples_ = samples; CvMat results_ = results; - if(RUN_PLAIN_IMPL) + if (RUN_PLAIN_IMPL) { CvSVM svm; svm.train(trainData, labels, Mat(), Mat(), params); @@ -197,7 +199,7 @@ PERF_TEST_P(SVMFixture, DISABLED_SVM, svm.predict(&samples_, &results_); } } - else if(RUN_OCL_IMPL) + else if (RUN_OCL_IMPL) { CvSVM_OCL svm; svm.train(trainData, labels, Mat(), Mat(), params); diff --git a/modules/ocl/perf/perf_norm.cpp b/modules/ocl/perf/perf_norm.cpp deleted file mode 100644 index abeb93cf1..000000000 --- a/modules/ocl/perf/perf_norm.cpp +++ /dev/null @@ -1,89 +0,0 @@ -/*M/////////////////////////////////////////////////////////////////////////////////////// -// -// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. -// -// By downloading, copying, installing or using the software you agree to this license. -// If you do not agree to this license, do not download, install, -// copy or use the software. -// -// -// License Agreement -// For Open Source Computer Vision Library -// -// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. -// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. -// Third party copyrights are property of their respective owners. -// -// @Authors -// Fangfang Bai, fangfang@multicorewareinc.com -// Jin Ma, jin@multicorewareinc.com -// -// 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 the copyright holders 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 "perf_precomp.hpp" - -using namespace perf; -using std::tr1::tuple; -using std::tr1::get; - -///////////// norm//////////////////////// - -CV_ENUM(NormType, NORM_INF, NORM_L1, NORM_L2) - -typedef std::tr1::tuple NormParams; -typedef TestBaseWithParam NormFixture; - -OCL_PERF_TEST_P(NormFixture, Norm, - ::testing::Combine(OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3), - OCL_TEST_TYPES, NormType::all())) -{ - const NormParams params = GetParam(); - const Size srcSize = get<0>(params); - const int type = get<1>(params); - const int normType = get<2>(params); - perf::ERROR_TYPE errorType = type != NORM_INF ? ERROR_RELATIVE : ERROR_ABSOLUTE; - double eps = 1e-5, value; - - Mat src1(srcSize, type), src2(srcSize, type); - declare.in(src1, src2, WARMUP_RNG); - - if (RUN_OCL_IMPL) - { - ocl::oclMat oclSrc1(src1), oclSrc2(src2); - - OCL_TEST_CYCLE() value = cv::ocl::norm(oclSrc1, oclSrc2, normType); - - SANITY_CHECK(value, eps, errorType); - } - else if (RUN_PLAIN_IMPL) - { - TEST_CYCLE() value = cv::norm(src1, src2, normType); - - SANITY_CHECK(value, eps, errorType); - } - else - OCL_PERF_ELSE -} diff --git a/modules/ocl/perf/perf_opticalflow.cpp b/modules/ocl/perf/perf_opticalflow.cpp index 3711e87d8..67079f392 100644 --- a/modules/ocl/perf/perf_opticalflow.cpp +++ b/modules/ocl/perf/perf_opticalflow.cpp @@ -52,13 +52,12 @@ using std::tr1::get; using std::tr1::tuple; using std::tr1::make_tuple; -typedef tuple PyrLKOpticalFlowParamType; -typedef TestBaseWithParam PyrLKOpticalFlowFixture; +typedef TestBaseWithParam > PyrLKOpticalFlowFixture; OCL_PERF_TEST_P(PyrLKOpticalFlowFixture, PyrLKOpticalFlow, ::testing::Values(1000, 2000, 4000)) { - const int pointsCount = GetParam(); + const int pointsCount = get<0>(GetParam()); const string fileName0 = "gpu/opticalflow/rubberwhale1.png", fileName1 = "gpu/opticalflow/rubberwhale2.png"; @@ -109,7 +108,7 @@ PERF_TEST(tvl1flowFixture, tvl1flow) const Size srcSize = frame0.size(); const double eps = 1.2; Mat flow(srcSize, CV_32FC2), flow1(srcSize, CV_32FC1), flow2(srcSize, CV_32FC1); - declare.in(frame0, frame1).out(flow1, flow2).time(159); + declare.in(frame0, frame1).out(flow1, flow2); if (RUN_PLAIN_IMPL) { diff --git a/modules/ocl/perf/perf_stat.cpp b/modules/ocl/perf/perf_stat.cpp new file mode 100644 index 000000000..a48cacb59 --- /dev/null +++ b/modules/ocl/perf/perf_stat.cpp @@ -0,0 +1,276 @@ +/*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. +// +// +// License Agreement +// For Open Source Computer Vision Library +// +// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Fangfang Bai, fangfang@multicorewareinc.com +// Jin Ma, jin@multicorewareinc.com +// +// 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 the copyright holders 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 "perf_precomp.hpp" + +using namespace perf; +using std::tr1::tuple; +using std::tr1::get; + + +///////////// MinMax //////////////////////// + +typedef Size_MatType MinMaxFixture; + +PERF_TEST_P(MinMaxFixture, MinMax, + ::testing::Combine(OCL_TYPICAL_MAT_SIZES, + OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +{ + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); + + Mat src(srcSize, type); + declare.in(src, WARMUP_RNG); + + double min_val = std::numeric_limits::max(), + max_val = std::numeric_limits::min(); + + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc(src); + + OCL_TEST_CYCLE() cv::ocl::minMax(oclSrc, &min_val, &max_val); + + ASSERT_GE(max_val, min_val); + SANITY_CHECK(min_val); + SANITY_CHECK(max_val); + } + else if (RUN_PLAIN_IMPL) + { + Point min_loc, max_loc; + + TEST_CYCLE() cv::minMaxLoc(src, &min_val, &max_val, &min_loc, &max_loc); + + ASSERT_GE(max_val, min_val); + SANITY_CHECK(min_val); + SANITY_CHECK(max_val); + } + else + OCL_PERF_ELSE +} + +///////////// MinMaxLoc //////////////////////// + +typedef Size_MatType MinMaxLocFixture; + +OCL_PERF_TEST_P(MinMaxLocFixture, MinMaxLoc, + ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +{ + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); + + Mat src(srcSize, type); + randu(src, 0, 1); + declare.in(src); + + double min_val = 0.0, max_val = 0.0; + Point min_loc, max_loc; + + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc(src); + + OCL_TEST_CYCLE() cv::ocl::minMaxLoc(oclSrc, &min_val, &max_val, &min_loc, &max_loc); + + ASSERT_GE(max_val, min_val); + SANITY_CHECK(min_val); + SANITY_CHECK(max_val); + } + else if (RUN_PLAIN_IMPL) + { + TEST_CYCLE() cv::minMaxLoc(src, &min_val, &max_val, &min_loc, &max_loc); + + ASSERT_GE(max_val, min_val); + SANITY_CHECK(min_val); + SANITY_CHECK(max_val); + } + else + OCL_PERF_ELSE +} + +///////////// Sum //////////////////////// + +typedef Size_MatType SumFixture; + +OCL_PERF_TEST_P(SumFixture, Sum, + ::testing::Combine(OCL_TEST_SIZES, + OCL_TEST_TYPES)) +{ + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); + + Mat src(srcSize, type); + Scalar result; + randu(src, 0, 60); + declare.in(src); + + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc(src); + + OCL_TEST_CYCLE() result = cv::ocl::sum(oclSrc); + + SANITY_CHECK(result, 1e-6, ERROR_RELATIVE); + } + else if (RUN_PLAIN_IMPL) + { + TEST_CYCLE() result = cv::sum(src); + + SANITY_CHECK(result, 1e-6, ERROR_RELATIVE); + } + else + OCL_PERF_ELSE +} + +///////////// countNonZero //////////////////////// + +typedef Size_MatType CountNonZeroFixture; + +OCL_PERF_TEST_P(CountNonZeroFixture, CountNonZero, + ::testing::Combine(OCL_TEST_SIZES, + OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +{ + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); + + Mat src(srcSize, type); + int result = 0; + randu(src, 0, 256); + declare.in(src); + + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc(src); + + OCL_TEST_CYCLE() result = cv::ocl::countNonZero(oclSrc); + + SANITY_CHECK(result); + } + else if (RUN_PLAIN_IMPL) + { + TEST_CYCLE() result = cv::countNonZero(src); + + SANITY_CHECK(result); + } + else + OCL_PERF_ELSE +} + +///////////// meanStdDev //////////////////////// + +typedef Size_MatType MeanStdDevFixture; + +OCL_PERF_TEST_P(MeanStdDevFixture, MeanStdDev, + ::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) +{ + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); + + Mat src(srcSize, type); + Scalar mean, stddev; + randu(src, 0, 256); + declare.in(src); + + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc(src); + + OCL_TEST_CYCLE() cv::ocl::meanStdDev(oclSrc, mean, stddev); + } + else if (RUN_PLAIN_IMPL) + { + TEST_CYCLE() cv::meanStdDev(src, mean, stddev); + } + else + OCL_PERF_ELSE + + SANITY_CHECK_NOTHING(); +// SANITY_CHECK(mean, 1e-6, ERROR_RELATIVE); +// SANITY_CHECK(stddev, 1e-6, ERROR_RELATIVE); +} + +///////////// norm//////////////////////// + +CV_ENUM(NormType, NORM_INF, NORM_L1, NORM_L2) + +typedef std::tr1::tuple NormParams; +typedef TestBaseWithParam NormFixture; + +OCL_PERF_TEST_P(NormFixture, Norm, + ::testing::Combine(OCL_PERF_ENUM(OCL_SIZE_1, OCL_SIZE_2, OCL_SIZE_3), + OCL_TEST_TYPES, NormType::all())) +{ + const NormParams params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params); + const int normType = get<2>(params); + perf::ERROR_TYPE errorType = type != NORM_INF ? ERROR_RELATIVE : ERROR_ABSOLUTE; + double eps = 1e-5, value; + + Mat src1(srcSize, type), src2(srcSize, type); + declare.in(src1, src2, WARMUP_RNG); + + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc1(src1), oclSrc2(src2); + + OCL_TEST_CYCLE() value = cv::ocl::norm(oclSrc1, oclSrc2, normType); + + SANITY_CHECK(value, eps, errorType); + } + else if (RUN_PLAIN_IMPL) + { + TEST_CYCLE() value = cv::norm(src1, src2, normType); + + SANITY_CHECK(value, eps, errorType); + } + else + OCL_PERF_ELSE +} From 29deff8707c825d1992a3bda6569a3e5eefb1a36 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sun, 2 Mar 2014 16:08:37 +0400 Subject: [PATCH 04/10] ability to merge logs with intersections only --- modules/ts/misc/summary.py | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/modules/ts/misc/summary.py b/modules/ts/misc/summary.py index 4e0cb7e7f..4cbb8901f 100755 --- a/modules/ts/misc/summary.py +++ b/modules/ts/misc/summary.py @@ -44,6 +44,7 @@ if __name__ == "__main__": parser.add_option("", "--match", dest="match", default=None) parser.add_option("", "--match-replace", dest="match_replace", default="") parser.add_option("", "--regressions-only", dest="regressionsOnly", default=None, metavar="X-FACTOR", help="show only tests with performance regressions not") + parser.add_option("", "--intersect-logs", dest="intersect_logs", default=False, help="show only tests present in all log files") (options, args) = parser.parse_args() options.generateHtml = detectHtmlOutputType(options.format) @@ -162,6 +163,10 @@ if __name__ == "__main__": for i in range(setsCount): case = cases[i] if case is None: + if options.intersect_logs: + needNewRow = False + break + tbl.newCell(str(i), "-") if options.calc_relatives and i > 0: tbl.newCell(str(i) + "%", "-") From 099ea91823f75ccab2976031140c1f25e1ebb8bb Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Mon, 3 Mar 2014 19:37:47 +0400 Subject: [PATCH 05/10] typos --- modules/calib3d/perf/perf_pnp.cpp | 2 +- modules/calib3d/test/test_solvepnp_ransac.cpp | 2 +- modules/contrib/src/imagelogpolprojection.cpp | 4 ++-- modules/contrib/src/lda.cpp | 2 +- modules/contrib/src/templatebuffer.hpp | 10 +++++----- modules/core/include/opencv2/core/core_c.h | 2 +- modules/core/src/array.cpp | 2 +- .../doc/feature_detection_and_description.rst | 6 +++--- .../include/opencv2/features2d/features2d.hpp | 4 ++-- modules/features2d/src/evaluation.cpp | 2 +- .../gpu/doc/feature_detection_and_description.rst | 12 ++++++------ modules/gpu/include/opencv2/gpu/gpu.hpp | 12 ++++++------ modules/gpu/perf/perf_features2d.cpp | 4 ++-- modules/gpu/perf/perf_imgproc.cpp | 6 +++--- modules/gpu/perf4au/main.cpp | 10 +++++----- modules/gpu/src/cuda/fast.cu | 8 ++++---- modules/gpu/src/fast.cpp | 14 +++++++------- modules/gpu/test/test_features2d.cpp | 14 +++++++------- modules/highgui/src/window_QT.cpp | 2 +- modules/highgui/src/window_carbon.cpp | 2 +- modules/imgproc/doc/feature_detection.rst | 2 +- .../imgproc/include/opencv2/imgproc/imgproc_c.h | 6 +++--- modules/imgproc/src/approx.cpp | 2 +- modules/imgproc/src/canny.cpp | 2 +- modules/ocl/src/opencl/haarobjectdetect.cl | 2 +- modules/ts/include/opencv2/ts/ts.hpp | 2 +- modules/viz/doc/widget.rst | 2 +- .../org/opencv/engine/manager/ManagerActivity.java | 14 +++++++------- samples/cpp/morphology2.cpp | 2 +- samples/gpu/morphology.cpp | 2 +- 30 files changed, 78 insertions(+), 78 deletions(-) diff --git a/modules/calib3d/perf/perf_pnp.cpp b/modules/calib3d/perf/perf_pnp.cpp index e0ffd70cf..557387d9a 100644 --- a/modules/calib3d/perf/perf_pnp.cpp +++ b/modules/calib3d/perf/perf_pnp.cpp @@ -126,7 +126,7 @@ PERF_TEST_P(PointsNum, DISABLED_SolvePnPRansac, testing::Values(4, 3*9, 7*13)) Mat tvec; #ifdef HAVE_TBB - // limit concurrency to get determenistic result + // limit concurrency to get deterministic result cv::Ptr one_thread = new tbb::task_scheduler_init(1); #endif diff --git a/modules/calib3d/test/test_solvepnp_ransac.cpp b/modules/calib3d/test/test_solvepnp_ransac.cpp index 6c924a580..4c8d56a96 100644 --- a/modules/calib3d/test/test_solvepnp_ransac.cpp +++ b/modules/calib3d/test/test_solvepnp_ransac.cpp @@ -271,7 +271,7 @@ TEST(DISABLED_Calib3d_SolvePnPRansac, concurrency) Mat tvec1, tvec2; { - // limit concurrency to get determenistic result + // limit concurrency to get deterministic result cv::theRNG().state = 20121010; cv::Ptr one_thread = new tbb::task_scheduler_init(1); solvePnPRansac(object, image, camera_mat, dist_coef, rvec1, tvec1); diff --git a/modules/contrib/src/imagelogpolprojection.cpp b/modules/contrib/src/imagelogpolprojection.cpp index ed821efa9..2824949c7 100644 --- a/modules/contrib/src/imagelogpolprojection.cpp +++ b/modules/contrib/src/imagelogpolprojection.cpp @@ -362,14 +362,14 @@ bool ImageLogPolProjection::_initLogPolarCortexSampling(const double reductionFa //std::cout<<"ImageLogPolProjection::Starting cortex projection"<size();++i) { - double curentValue=(double)*(bufferPTR++); + double currentValue=(double)*(bufferPTR++); // updating "closest to the high threshold" pixel value - double highValueTest=maxThreshold-curentValue; + double highValueTest=maxThreshold-currentValue; if (highValueTest>0) { if (deltaH>highValueTest) { deltaH=highValueTest; - updatedHighValue=curentValue; + updatedHighValue=currentValue; } } // updating "closest to the low threshold" pixel value - double lowValueTest=curentValue-minThreshold; + double lowValueTest=currentValue-minThreshold; if (lowValueTest>0) { if (deltaL>lowValueTest) { deltaL=lowValueTest; - updatedLowValue=curentValue; + updatedLowValue=currentValue; } } } diff --git a/modules/core/include/opencv2/core/core_c.h b/modules/core/include/opencv2/core/core_c.h index d4182d2f7..38abfc409 100644 --- a/modules/core/include/opencv2/core/core_c.h +++ b/modules/core/include/opencv2/core/core_c.h @@ -105,7 +105,7 @@ CVAPI(void) cvResetImageROI( IplImage* image ); /* Retrieves image ROI */ CVAPI(CvRect) cvGetImageROI( const IplImage* image ); -/* Allocates and initalizes CvMat header */ +/* Allocates and initializes CvMat header */ CVAPI(CvMat*) cvCreateMatHeader( int rows, int cols, int type ); #define CV_AUTOSTEP 0x7fffffff diff --git a/modules/core/src/array.cpp b/modules/core/src/array.cpp index 9c024293e..a35ad0e2a 100644 --- a/modules/core/src/array.cpp +++ b/modules/core/src/array.cpp @@ -2897,7 +2897,7 @@ cvCreateImage( CvSize size, int depth, int channels ) } -// initalize IplImage header, allocated by the user +// initialize IplImage header, allocated by the user CV_IMPL IplImage* cvInitImageHeader( IplImage * image, CvSize size, int depth, int channels, int origin, int align ) diff --git a/modules/features2d/doc/feature_detection_and_description.rst b/modules/features2d/doc/feature_detection_and_description.rst index a0272026d..c05a71d34 100644 --- a/modules/features2d/doc/feature_detection_and_description.rst +++ b/modules/features2d/doc/feature_detection_and_description.rst @@ -11,9 +11,9 @@ FAST ---- Detects corners using the FAST algorithm -.. ocv:function:: void FAST( InputArray image, vector& keypoints, int threshold, bool nonmaxSupression=true ) +.. ocv:function:: void FAST( InputArray image, vector& keypoints, int threshold, bool nonmaxSuppression=true ) -.. ocv:function:: void FASTX( InputArray image, vector& keypoints, int threshold, bool nonmaxSupression, int type ) +.. ocv:function:: void FASTX( InputArray image, vector& keypoints, int threshold, bool nonmaxSuppression, int type ) :param image: grayscale image where keypoints (corners) are detected. @@ -21,7 +21,7 @@ Detects corners using the FAST algorithm :param threshold: threshold on difference between intensity of the central pixel and pixels of a circle around this pixel. - :param nonmaxSupression: if true, non-maximum suppression is applied to detected corners (keypoints). + :param nonmaxSuppression: if true, non-maximum suppression is applied to detected corners (keypoints). :param type: one of the three neighborhoods as defined in the paper: ``FastFeatureDetector::TYPE_9_16``, ``FastFeatureDetector::TYPE_7_12``, ``FastFeatureDetector::TYPE_5_8`` diff --git a/modules/features2d/include/opencv2/features2d/features2d.hpp b/modules/features2d/include/opencv2/features2d/features2d.hpp index d598ed6d7..7536128c7 100644 --- a/modules/features2d/include/opencv2/features2d/features2d.hpp +++ b/modules/features2d/include/opencv2/features2d/features2d.hpp @@ -566,10 +566,10 @@ protected: //! detects corners using FAST algorithm by E. Rosten CV_EXPORTS void FAST( InputArray image, CV_OUT vector& keypoints, - int threshold, bool nonmaxSupression=true ); + int threshold, bool nonmaxSuppression=true ); CV_EXPORTS void FASTX( InputArray image, CV_OUT vector& keypoints, - int threshold, bool nonmaxSupression, int type ); + int threshold, bool nonmaxSuppression, int type ); class CV_EXPORTS_W FastFeatureDetector : public FeatureDetector { diff --git a/modules/features2d/src/evaluation.cpp b/modules/features2d/src/evaluation.cpp index 44151c03c..52740e6e5 100644 --- a/modules/features2d/src/evaluation.cpp +++ b/modules/features2d/src/evaluation.cpp @@ -128,7 +128,7 @@ public: Point2f center; Scalar ellipse; // 3 elements a, b, c: ax^2+2bxy+cy^2=1 - Size_ axes; // half lenght of elipse axes + Size_ axes; // half length of ellipse axes Size_ boundingBox; // half sizes of bounding box which sides are parallel to the coordinate axes }; diff --git a/modules/gpu/doc/feature_detection_and_description.rst b/modules/gpu/doc/feature_detection_and_description.rst index 191b0ada0..66f425004 100644 --- a/modules/gpu/doc/feature_detection_and_description.rst +++ b/modules/gpu/doc/feature_detection_and_description.rst @@ -24,7 +24,7 @@ Class used for corner detection using the FAST algorithm. :: // all features have same size static const int FEATURE_SIZE = 7; - explicit FAST_GPU(int threshold, bool nonmaxSupression = true, + explicit FAST_GPU(int threshold, bool nonmaxSuppression = true, double keypointsRatio = 0.05); void operator ()(const GpuMat& image, const GpuMat& mask, GpuMat& keypoints); @@ -39,7 +39,7 @@ Class used for corner detection using the FAST algorithm. :: void release(); - bool nonmaxSupression; + bool nonmaxSuppression; int threshold; @@ -61,11 +61,11 @@ gpu::FAST_GPU::FAST_GPU ------------------------------------- Constructor. -.. ocv:function:: gpu::FAST_GPU::FAST_GPU(int threshold, bool nonmaxSupression = true, double keypointsRatio = 0.05) +.. ocv:function:: gpu::FAST_GPU::FAST_GPU(int threshold, bool nonmaxSuppression = true, double keypointsRatio = 0.05) :param threshold: Threshold on difference between intensity of the central pixel and pixels on a circle around this pixel. - :param nonmaxSupression: If it is true, non-maximum suppression is applied to detected corners (keypoints). + :param nonmaxSuppression: If it is true, non-maximum suppression is applied to detected corners (keypoints). :param keypointsRatio: Inner buffer size for keypoints store is determined as (keypointsRatio * image_width * image_height). @@ -115,7 +115,7 @@ Releases inner buffer memory. gpu::FAST_GPU::calcKeyPointsLocation ------------------------------------- -Find keypoints and compute it's response if ``nonmaxSupression`` is true. +Find keypoints and compute it's response if ``nonmaxSuppression`` is true. .. ocv:function:: int gpu::FAST_GPU::calcKeyPointsLocation(const GpuMat& image, const GpuMat& mask) @@ -185,7 +185,7 @@ Class for extracting ORB features and descriptors from an image. :: int descriptorSize() const; void setParams(size_t n_features, const ORB::CommonParams& detector_params); - void setFastParams(int threshold, bool nonmaxSupression = true); + void setFastParams(int threshold, bool nonmaxSuppression = true); void release(); diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 053bcdbc5..0ab0fb1cd 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1580,7 +1580,7 @@ public: // all features have same size static const int FEATURE_SIZE = 7; - explicit FAST_GPU(int threshold, bool nonmaxSupression = true, double keypointsRatio = 0.05); + explicit FAST_GPU(int threshold, bool nonmaxSuppression = true, double keypointsRatio = 0.05); //! finds the keypoints using FAST detector //! supports only CV_8UC1 images @@ -1596,19 +1596,19 @@ public: //! release temporary buffer's memory void release(); - bool nonmaxSupression; + bool nonmaxSuppression; int threshold; //! max keypoints = keypointsRatio * img.size().area() double keypointsRatio; - //! find keypoints and compute it's response if nonmaxSupression is true + //! find keypoints and compute it's response if nonmaxSuppression is true //! return count of detected keypoints int calcKeyPointsLocation(const GpuMat& image, const GpuMat& mask); //! get final array of keypoints - //! performs nonmax supression if needed + //! performs nonmax suppression if needed //! return final count of keypoints int getKeyPoints(GpuMat& keypoints); @@ -1670,10 +1670,10 @@ public: //! returns the descriptor size in bytes inline int descriptorSize() const { return kBytes; } - inline void setFastParams(int threshold, bool nonmaxSupression = true) + inline void setFastParams(int threshold, bool nonmaxSuppression = true) { fastDetector_.threshold = threshold; - fastDetector_.nonmaxSupression = nonmaxSupression; + fastDetector_.nonmaxSuppression = nonmaxSuppression; } //! release temporary buffer's memory diff --git a/modules/gpu/perf/perf_features2d.cpp b/modules/gpu/perf/perf_features2d.cpp index feee3a939..45823cef2 100644 --- a/modules/gpu/perf/perf_features2d.cpp +++ b/modules/gpu/perf/perf_features2d.cpp @@ -49,9 +49,9 @@ using namespace perf; ////////////////////////////////////////////////////////////////////// // FAST -DEF_PARAM_TEST(Image_Threshold_NonMaxSupression, string, int, bool); +DEF_PARAM_TEST(Image_Threshold_NonMaxSuppression, string, int, bool); -PERF_TEST_P(Image_Threshold_NonMaxSupression, Features2D_FAST, +PERF_TEST_P(Image_Threshold_NonMaxSuppression, Features2D_FAST, Combine(Values("gpu/perf/aloe.png"), Values(20), Bool())) diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp index be0e312a0..fa0a42cfe 100644 --- a/modules/gpu/perf/perf_imgproc.cpp +++ b/modules/gpu/perf/perf_imgproc.cpp @@ -1746,7 +1746,7 @@ PERF_TEST_P(Image, ImgProc_HoughLinesP, const float rho = 1.0f; const float theta = static_cast(CV_PI / 180.0); const int threshold = 100; - const int minLineLenght = 50; + const int minLineLength = 50; const int maxLineGap = 5; const cv::Mat image = cv::imread(fileName, cv::IMREAD_GRAYSCALE); @@ -1761,7 +1761,7 @@ PERF_TEST_P(Image, ImgProc_HoughLinesP, cv::gpu::GpuMat d_lines; cv::gpu::HoughLinesBuf d_buf; - TEST_CYCLE() cv::gpu::HoughLinesP(d_mask, d_lines, d_buf, rho, theta, minLineLenght, maxLineGap); + TEST_CYCLE() cv::gpu::HoughLinesP(d_mask, d_lines, d_buf, rho, theta, minLineLength, maxLineGap); cv::Mat gpu_lines(d_lines); cv::Vec4i* begin = gpu_lines.ptr(); @@ -1773,7 +1773,7 @@ PERF_TEST_P(Image, ImgProc_HoughLinesP, { std::vector cpu_lines; - TEST_CYCLE() cv::HoughLinesP(mask, cpu_lines, rho, theta, threshold, minLineLenght, maxLineGap); + TEST_CYCLE() cv::HoughLinesP(mask, cpu_lines, rho, theta, threshold, minLineLength, maxLineGap); SANITY_CHECK(cpu_lines); } diff --git a/modules/gpu/perf4au/main.cpp b/modules/gpu/perf4au/main.cpp index 8c68385e1..ceee61e1d 100644 --- a/modules/gpu/perf4au/main.cpp +++ b/modules/gpu/perf4au/main.cpp @@ -74,7 +74,7 @@ PERF_TEST_P(Image, HoughLinesP, testing::Values(std::string("im1_1280x800.jpg")) const float rho = 1.f; const float theta = 1.f; const int threshold = 40; - const int minLineLenght = 20; + const int minLineLength = 20; const int maxLineGap = 5; cv::Mat image = cv::imread(fileName, cv::IMREAD_GRAYSCALE); @@ -85,11 +85,11 @@ PERF_TEST_P(Image, HoughLinesP, testing::Values(std::string("im1_1280x800.jpg")) cv::gpu::GpuMat d_lines; cv::gpu::HoughLinesBuf d_buf; - cv::gpu::HoughLinesP(d_image, d_lines, d_buf, rho, theta, minLineLenght, maxLineGap); + cv::gpu::HoughLinesP(d_image, d_lines, d_buf, rho, theta, minLineLength, maxLineGap); TEST_CYCLE() { - cv::gpu::HoughLinesP(d_image, d_lines, d_buf, rho, theta, minLineLenght, maxLineGap); + cv::gpu::HoughLinesP(d_image, d_lines, d_buf, rho, theta, minLineLength, maxLineGap); } } else @@ -98,11 +98,11 @@ PERF_TEST_P(Image, HoughLinesP, testing::Values(std::string("im1_1280x800.jpg")) cv::Canny(image, mask, 50, 100); std::vector lines; - cv::HoughLinesP(mask, lines, rho, theta, threshold, minLineLenght, maxLineGap); + cv::HoughLinesP(mask, lines, rho, theta, threshold, minLineLength, maxLineGap); TEST_CYCLE() { - cv::HoughLinesP(mask, lines, rho, theta, threshold, minLineLenght, maxLineGap); + cv::HoughLinesP(mask, lines, rho, theta, threshold, minLineLength, maxLineGap); } } diff --git a/modules/gpu/src/cuda/fast.cu b/modules/gpu/src/cuda/fast.cu index 9bf1f7ad9..f72b1fc04 100644 --- a/modules/gpu/src/cuda/fast.cu +++ b/modules/gpu/src/cuda/fast.cu @@ -318,9 +318,9 @@ namespace cv { namespace gpu { namespace device } /////////////////////////////////////////////////////////////////////////// - // nonmaxSupression + // nonmaxSuppression - __global__ void nonmaxSupression(const short2* kpLoc, int count, const PtrStepSzi scoreMat, short2* locFinal, float* responseFinal) + __global__ void nonmaxSuppression(const short2* kpLoc, int count, const PtrStepSzi scoreMat, short2* locFinal, float* responseFinal) { #if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 110) @@ -356,7 +356,7 @@ namespace cv { namespace gpu { namespace device #endif } - int nonmaxSupression_gpu(const short2* kpLoc, int count, PtrStepSzi score, short2* loc, float* response) + int nonmaxSuppression_gpu(const short2* kpLoc, int count, PtrStepSzi score, short2* loc, float* response) { void* counter_ptr; cudaSafeCall( cudaGetSymbolAddress(&counter_ptr, g_counter) ); @@ -368,7 +368,7 @@ namespace cv { namespace gpu { namespace device cudaSafeCall( cudaMemset(counter_ptr, 0, sizeof(unsigned int)) ); - nonmaxSupression<<>>(kpLoc, count, score, loc, response); + nonmaxSuppression<<>>(kpLoc, count, score, loc, response); cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaDeviceSynchronize() ); diff --git a/modules/gpu/src/fast.cpp b/modules/gpu/src/fast.cpp index a24a69962..37ca92dc4 100644 --- a/modules/gpu/src/fast.cpp +++ b/modules/gpu/src/fast.cpp @@ -59,8 +59,8 @@ int cv::gpu::FAST_GPU::getKeyPoints(GpuMat&) { throw_nogpu(); return 0; } #else /* !defined (HAVE_CUDA) */ -cv::gpu::FAST_GPU::FAST_GPU(int _threshold, bool _nonmaxSupression, double _keypointsRatio) : - nonmaxSupression(_nonmaxSupression), threshold(_threshold), keypointsRatio(_keypointsRatio), count_(0) +cv::gpu::FAST_GPU::FAST_GPU(int _threshold, bool _nonmaxSuppression, double _keypointsRatio) : + nonmaxSuppression(_nonmaxSuppression), threshold(_threshold), keypointsRatio(_keypointsRatio), count_(0) { } @@ -114,7 +114,7 @@ namespace cv { namespace gpu { namespace device namespace fast { int calcKeypoints_gpu(PtrStepSzb img, PtrStepSzb mask, short2* kpLoc, int maxKeypoints, PtrStepSzi score, int threshold); - int nonmaxSupression_gpu(const short2* kpLoc, int count, PtrStepSzi score, short2* loc, float* response); + int nonmaxSuppression_gpu(const short2* kpLoc, int count, PtrStepSzi score, short2* loc, float* response); } }}} @@ -129,13 +129,13 @@ int cv::gpu::FAST_GPU::calcKeyPointsLocation(const GpuMat& img, const GpuMat& ma ensureSizeIsEnough(1, maxKeypoints, CV_16SC2, kpLoc_); - if (nonmaxSupression) + if (nonmaxSuppression) { ensureSizeIsEnough(img.size(), CV_32SC1, score_); score_.setTo(Scalar::all(0)); } - count_ = calcKeypoints_gpu(img, mask, kpLoc_.ptr(), maxKeypoints, nonmaxSupression ? score_ : PtrStepSzi(), threshold); + count_ = calcKeypoints_gpu(img, mask, kpLoc_.ptr(), maxKeypoints, nonmaxSuppression ? score_ : PtrStepSzi(), threshold); count_ = std::min(count_, maxKeypoints); return count_; @@ -150,8 +150,8 @@ int cv::gpu::FAST_GPU::getKeyPoints(GpuMat& keypoints) ensureSizeIsEnough(ROWS_COUNT, count_, CV_32FC1, keypoints); - if (nonmaxSupression) - return nonmaxSupression_gpu(kpLoc_.ptr(), count_, score_, keypoints.ptr(LOCATION_ROW), keypoints.ptr(RESPONSE_ROW)); + if (nonmaxSuppression) + return nonmaxSuppression_gpu(kpLoc_.ptr(), count_, score_, keypoints.ptr(LOCATION_ROW), keypoints.ptr(RESPONSE_ROW)); GpuMat locRow(1, count_, kpLoc_.type(), keypoints.ptr(0)); kpLoc_.colRange(0, count_).copyTo(locRow); diff --git a/modules/gpu/test/test_features2d.cpp b/modules/gpu/test/test_features2d.cpp index dfa3afa38..697483657 100644 --- a/modules/gpu/test/test_features2d.cpp +++ b/modules/gpu/test/test_features2d.cpp @@ -52,20 +52,20 @@ using namespace cvtest; namespace { IMPLEMENT_PARAM_CLASS(FAST_Threshold, int) - IMPLEMENT_PARAM_CLASS(FAST_NonmaxSupression, bool) + IMPLEMENT_PARAM_CLASS(FAST_NonmaxSuppression, bool) } -PARAM_TEST_CASE(FAST, cv::gpu::DeviceInfo, FAST_Threshold, FAST_NonmaxSupression) +PARAM_TEST_CASE(FAST, cv::gpu::DeviceInfo, FAST_Threshold, FAST_NonmaxSuppression) { cv::gpu::DeviceInfo devInfo; int threshold; - bool nonmaxSupression; + bool nonmaxSuppression; virtual void SetUp() { devInfo = GET_PARAM(0); threshold = GET_PARAM(1); - nonmaxSupression = GET_PARAM(2); + nonmaxSuppression = GET_PARAM(2); cv::gpu::setDevice(devInfo.deviceID()); } @@ -77,7 +77,7 @@ GPU_TEST_P(FAST, Accuracy) ASSERT_FALSE(image.empty()); cv::gpu::FAST_GPU fast(threshold); - fast.nonmaxSupression = nonmaxSupression; + fast.nonmaxSuppression = nonmaxSuppression; if (!supportFeature(devInfo, cv::gpu::GLOBAL_ATOMICS)) { @@ -97,7 +97,7 @@ GPU_TEST_P(FAST, Accuracy) fast(loadMat(image), cv::gpu::GpuMat(), keypoints); std::vector keypoints_gold; - cv::FAST(image, keypoints_gold, threshold, nonmaxSupression); + cv::FAST(image, keypoints_gold, threshold, nonmaxSuppression); ASSERT_KEYPOINTS_EQ(keypoints_gold, keypoints); } @@ -106,7 +106,7 @@ GPU_TEST_P(FAST, Accuracy) INSTANTIATE_TEST_CASE_P(GPU_Features2D, FAST, testing::Combine( ALL_DEVICES, testing::Values(FAST_Threshold(25), FAST_Threshold(50)), - testing::Values(FAST_NonmaxSupression(false), FAST_NonmaxSupression(true)))); + testing::Values(FAST_NonmaxSuppression(false), FAST_NonmaxSuppression(true)))); ///////////////////////////////////////////////////////////////////////////////////////////////// // ORB diff --git a/modules/highgui/src/window_QT.cpp b/modules/highgui/src/window_QT.cpp index 84af3941a..64ebf0870 100644 --- a/modules/highgui/src/window_QT.cpp +++ b/modules/highgui/src/window_QT.cpp @@ -1536,7 +1536,7 @@ CvWindow::CvWindow(QString name, int arg2) setWindowTitle(name); setObjectName(name); - setFocus( Qt::PopupFocusReason ); //#1695 arrow keys are not recieved without the explicit focus + setFocus( Qt::PopupFocusReason ); //#1695 arrow keys are not received without the explicit focus resize(400, 300); setMinimumSize(1, 1); diff --git a/modules/highgui/src/window_carbon.cpp b/modules/highgui/src/window_carbon.cpp index 722df2cdd..3d092e736 100644 --- a/modules/highgui/src/window_carbon.cpp +++ b/modules/highgui/src/window_carbon.cpp @@ -569,7 +569,7 @@ static int icvCreateTrackbar (const char* trackbar_name, //pad size maxvalue in pixel Point qdSize; - char valueinchar[strlen(trackbar_name)+1 +1 +1+nbDigit+1];//lenght+\n +space +(+nbDigit+) + char valueinchar[strlen(trackbar_name)+1 +1 +1+nbDigit+1];//length+\n +space +(+nbDigit+) sprintf(valueinchar, "%s (%d)",trackbar_name, trackbar->maxval); SInt16 baseline; CFStringRef text = CFStringCreateWithCString(NULL,valueinchar,kCFStringEncodingASCII); diff --git a/modules/imgproc/doc/feature_detection.rst b/modules/imgproc/doc/feature_detection.rst index 4f922f2a7..861c16432 100644 --- a/modules/imgproc/doc/feature_detection.rst +++ b/modules/imgproc/doc/feature_detection.rst @@ -101,7 +101,7 @@ Harris edge detector. .. ocv:pyfunction:: cv2.cornerHarris(src, blockSize, ksize, k[, dst[, borderType]]) -> dst -.. ocv:cfunction:: void cvCornerHarris( const CvArr* image, CvArr* harris_responce, int block_size, int aperture_size=3, double k=0.04 ) +.. ocv:cfunction:: void cvCornerHarris( const CvArr* image, CvArr* harris_response, int block_size, int aperture_size=3, double k=0.04 ) .. ocv:pyoldfunction:: cv.CornerHarris(image, harris_dst, blockSize, aperture_size=3, k=0.04) -> None diff --git a/modules/imgproc/include/opencv2/imgproc/imgproc_c.h b/modules/imgproc/include/opencv2/imgproc/imgproc_c.h index 4ba1b2b26..46d9f0139 100644 --- a/modules/imgproc/include/opencv2/imgproc/imgproc_c.h +++ b/modules/imgproc/include/opencv2/imgproc/imgproc_c.h @@ -304,7 +304,7 @@ CVAPI(int) cvFindContours( CvArr* image, CvMemStorage* storage, CvSeq** first_c int method CV_DEFAULT(CV_CHAIN_APPROX_SIMPLE), CvPoint offset CV_DEFAULT(cvPoint(0,0))); -/* Initalizes contour retrieving process. +/* Initializes contour retrieving process. Calls cvStartFindContours. Calls cvFindNextContour until null pointer is returned or some other condition becomes true. @@ -334,7 +334,7 @@ CVAPI(CvSeq*) cvApproxChains( CvSeq* src_seq, CvMemStorage* storage, int minimal_perimeter CV_DEFAULT(0), int recursive CV_DEFAULT(0)); -/* Initalizes Freeman chain reader. +/* Initializes Freeman chain reader. The reader is used to iteratively get coordinates of all the chain points. If the Freeman codes should be read as is, a simple sequence reader should be used */ CVAPI(void) cvStartReadChainPoints( CvChain* chain, CvChainPtReader* reader ); @@ -573,7 +573,7 @@ CVAPI(void) cvCornerMinEigenVal( const CvArr* image, CvArr* eigenval, /* Harris corner detector: Calculates det(M) - k*(trace(M)^2), where M is 2x2 gradient covariation matrix for each pixel */ -CVAPI(void) cvCornerHarris( const CvArr* image, CvArr* harris_responce, +CVAPI(void) cvCornerHarris( const CvArr* image, CvArr* harris_response, int block_size, int aperture_size CV_DEFAULT(3), double k CV_DEFAULT(0.04) ); diff --git a/modules/imgproc/src/approx.cpp b/modules/imgproc/src/approx.cpp index c2831ca70..8bd76c3cb 100644 --- a/modules/imgproc/src/approx.cpp +++ b/modules/imgproc/src/approx.cpp @@ -220,7 +220,7 @@ CvSeq* icvApproximateChainTC89( CvChain* chain, int header_size, current = temp.next; /* Pass 2. - Performs non-maxima supression */ + Performs non-maxima suppression */ do { int k2 = current->k >> 1; diff --git a/modules/imgproc/src/canny.cpp b/modules/imgproc/src/canny.cpp index 44fd42a2a..2161db91c 100644 --- a/modules/imgproc/src/canny.cpp +++ b/modules/imgproc/src/canny.cpp @@ -171,7 +171,7 @@ void cv::Canny( InputArray _src, OutputArray _dst, #define CANNY_PUSH(d) *(d) = uchar(2), *stack_top++ = (d) #define CANNY_POP(d) (d) = *--stack_top - // calculate magnitude and angle of gradient, perform non-maxima supression. + // calculate magnitude and angle of gradient, perform non-maxima suppression. // fill the map with one of the following values: // 0 - the pixel might belong to an edge // 1 - the pixel can not belong to an edge diff --git a/modules/ocl/src/opencl/haarobjectdetect.cl b/modules/ocl/src/opencl/haarobjectdetect.cl index 8464a580b..39d11b0e7 100644 --- a/modules/ocl/src/opencl/haarobjectdetect.cl +++ b/modules/ocl/src/opencl/haarobjectdetect.cl @@ -217,7 +217,7 @@ __kernel void gpuRunHaarClassifierCascadePacked( (SumL[M0(n0.x)+lcl_off] - SumL[M1(n0.x)+lcl_off] - SumL[M0(n0.y)+lcl_off] + SumL[M1(n0.y)+lcl_off]) * as_float(n1.z) + (SumL[M0(n0.z)+lcl_off] - SumL[M1(n0.z)+lcl_off] - SumL[M0(n0.w)+lcl_off] + SumL[M1(n0.w)+lcl_off]) * as_float(n1.w) + (SumL[M0(n1.x)+lcl_off] - SumL[M1(n1.x)+lcl_off] - SumL[M0(n1.y)+lcl_off] + SumL[M1(n1.y)+lcl_off]) * as_float(n2.x); - //accumulate stage responce + //accumulate stage response stage_sum += (classsum >= nodethreshold) ? as_float(n2.w) : as_float(n2.z); } result = (stage_sum >= stagethreshold); diff --git a/modules/ts/include/opencv2/ts/ts.hpp b/modules/ts/include/opencv2/ts/ts.hpp index 8ea1ad93b..f4cd3ffbc 100644 --- a/modules/ts/include/opencv2/ts/ts.hpp +++ b/modules/ts/include/opencv2/ts/ts.hpp @@ -372,7 +372,7 @@ public: // processing time (in this case there should be possibility to interrupt such a function FAIL_HANG=-13, - // unexpected responce on passing bad arguments to the tested function + // unexpected response on passing bad arguments to the tested function // (the function crashed, proceed succesfully (while it should not), or returned // error code that is different from what is expected) FAIL_BAD_ARG_CHECK=-14, diff --git a/modules/viz/doc/widget.rst b/modules/viz/doc/widget.rst index 008e0e68a..7f8926d3f 100644 --- a/modules/viz/doc/widget.rst +++ b/modules/viz/doc/widget.rst @@ -397,7 +397,7 @@ This 3D Widget defines a cone. :: { public: //! create default cone, oriented along x-axis with center of its base located at origin - WCone(double lenght, double radius, int resolution = 6.0, const Color &color = Color::white()); + WCone(double length, double radius, int resolution = 6.0, const Color &color = Color::white()); //! creates repositioned cone WCone(double radius, const Point3d& center, const Point3d& tip, int resolution = 6.0, const Color &color = Color::white()); diff --git a/platforms/android/service/engine/src/org/opencv/engine/manager/ManagerActivity.java b/platforms/android/service/engine/src/org/opencv/engine/manager/ManagerActivity.java index 8e8389dcc..8d4024203 100644 --- a/platforms/android/service/engine/src/org/opencv/engine/manager/ManagerActivity.java +++ b/platforms/android/service/engine/src/org/opencv/engine/manager/ManagerActivity.java @@ -214,12 +214,12 @@ public class ManagerActivity extends Activity } }); - mPackageChangeReciever = new BroadcastReceiver() { + mPackageChangeReceiver = new BroadcastReceiver() { @Override public void onReceive(Context context, Intent intent) { - Log.d("OpenCVManager/Reciever", "Bradcast message " + intent.getAction() + " reciever"); - Log.d("OpenCVManager/Reciever", "Filling package list on broadcast message"); + Log.d("OpenCVManager/Receiver", "Broadcast message " + intent.getAction() + " receiver"); + Log.d("OpenCVManager/Receiver", "Filling package list on broadcast message"); if (!bindService(new Intent("org.opencv.engine.BIND"), new OpenCVEngineServiceConnection(), Context.BIND_AUTO_CREATE)) { TextView EngineVersionView = (TextView)findViewById(R.id.EngineVersionValue); @@ -235,14 +235,14 @@ public class ManagerActivity extends Activity filter.addAction(Intent.ACTION_PACKAGE_REMOVED); filter.addAction(Intent.ACTION_PACKAGE_REPLACED); - registerReceiver(mPackageChangeReciever, filter); + registerReceiver(mPackageChangeReceiver, filter); } @Override protected void onDestroy() { super.onDestroy(); - if (mPackageChangeReciever != null) - unregisterReceiver(mPackageChangeReciever); + if (mPackageChangeReceiver != null) + unregisterReceiver(mPackageChangeReceiver); } @Override @@ -273,7 +273,7 @@ public class ManagerActivity extends Activity protected int ManagerApiLevel = 0; protected String ManagerVersion; - protected BroadcastReceiver mPackageChangeReciever = null; + protected BroadcastReceiver mPackageChangeReceiver = null; protected class OpenCVEngineServiceConnection implements ServiceConnection { diff --git a/samples/cpp/morphology2.cpp b/samples/cpp/morphology2.cpp index 330f4e064..a7793008a 100644 --- a/samples/cpp/morphology2.cpp +++ b/samples/cpp/morphology2.cpp @@ -12,7 +12,7 @@ static void help() printf("\nShow off image morphology: erosion, dialation, open and close\n" "Call:\n morphology2 [image]\n" - "This program also shows use of rect, elipse and cross kernels\n\n"); + "This program also shows use of rect, ellipse and cross kernels\n\n"); printf( "Hot keys: \n" "\tESC - quit the program\n" "\tr - use rectangle structuring element\n" diff --git a/samples/gpu/morphology.cpp b/samples/gpu/morphology.cpp index 36f518b70..2596b2d0f 100644 --- a/samples/gpu/morphology.cpp +++ b/samples/gpu/morphology.cpp @@ -13,7 +13,7 @@ static void help() printf("\nShow off image morphology: erosion, dialation, open and close\n" "Call:\n morphology2 [image]\n" - "This program also shows use of rect, elipse and cross kernels\n\n"); + "This program also shows use of rect, ellipse and cross kernels\n\n"); printf( "Hot keys: \n" "\tESC - quit the program\n" "\tr - use rectangle structuring element\n" From b184d7f27b9115b7532c7e021e11051a74e93fde Mon Sep 17 00:00:00 2001 From: Dave Hughes Date: Wed, 5 Mar 2014 02:44:40 +0000 Subject: [PATCH 06/10] Fix for #3554 v4l2_scan_controls_enumerate_menu is unused and causes ioctl error on RaspberryPi and possibly other Video4Linux variants. See http://www.raspberrypi.org/forum/viewtopic.php?f=43&t=65026 for more detail. --- modules/highgui/src/cap_libv4l.cpp | 29 ----------------------------- modules/highgui/src/cap_v4l.cpp | 24 ------------------------ 2 files changed, 53 deletions(-) diff --git a/modules/highgui/src/cap_libv4l.cpp b/modules/highgui/src/cap_libv4l.cpp index 91047de1f..e7aa5b5df 100644 --- a/modules/highgui/src/cap_libv4l.cpp +++ b/modules/highgui/src/cap_libv4l.cpp @@ -321,7 +321,6 @@ typedef struct CvCaptureCAM_V4L struct v4l2_control control; enum v4l2_buf_type type; struct v4l2_queryctrl queryctrl; - struct v4l2_querymenu querymenu; /* V4L2 control variables */ v4l2_ctrl_range** v4l2_ctrl_ranges; @@ -491,25 +490,6 @@ static int try_init_v4l2(CvCaptureCAM_V4L* capture, char *deviceName) } -static void v4l2_scan_controls_enumerate_menu(CvCaptureCAM_V4L* capture) -{ -// printf (" Menu items:\n"); - CLEAR (capture->querymenu); - capture->querymenu.id = capture->queryctrl.id; - for (capture->querymenu.index = capture->queryctrl.minimum; - (int)capture->querymenu.index <= capture->queryctrl.maximum; - capture->querymenu.index++) - { - if (0 == xioctl (capture->deviceHandle, VIDIOC_QUERYMENU, - &capture->querymenu)) - { - //printf (" %s\n", capture->querymenu.name); - } else { - perror ("VIDIOC_QUERYMENU"); - } - } -} - static void v4l2_free_ranges(CvCaptureCAM_V4L* capture) { int i; if (capture->v4l2_ctrl_ranges != NULL) { @@ -590,9 +570,6 @@ static void v4l2_scan_controls(CvCaptureCAM_V4L* capture) { if(capture->queryctrl.flags & V4L2_CTRL_FLAG_DISABLED) { continue; } - if (capture->queryctrl.type == V4L2_CTRL_TYPE_MENU) { - v4l2_scan_controls_enumerate_menu(capture); - } if(capture->queryctrl.type != V4L2_CTRL_TYPE_INTEGER && capture->queryctrl.type != V4L2_CTRL_TYPE_BOOLEAN && capture->queryctrl.type != V4L2_CTRL_TYPE_MENU) { @@ -613,9 +590,6 @@ static void v4l2_scan_controls(CvCaptureCAM_V4L* capture) { if(capture->queryctrl.flags & V4L2_CTRL_FLAG_DISABLED) { continue; } - if (capture->queryctrl.type == V4L2_CTRL_TYPE_MENU) { - v4l2_scan_controls_enumerate_menu(capture); - } if(capture->queryctrl.type != V4L2_CTRL_TYPE_INTEGER && capture->queryctrl.type != V4L2_CTRL_TYPE_BOOLEAN && capture->queryctrl.type != V4L2_CTRL_TYPE_MENU) { @@ -637,9 +611,6 @@ static void v4l2_scan_controls(CvCaptureCAM_V4L* capture) { continue; } - if (capture->queryctrl.type == V4L2_CTRL_TYPE_MENU) { - v4l2_scan_controls_enumerate_menu(capture); - } if(capture->queryctrl.type != V4L2_CTRL_TYPE_INTEGER && capture->queryctrl.type != V4L2_CTRL_TYPE_BOOLEAN && diff --git a/modules/highgui/src/cap_v4l.cpp b/modules/highgui/src/cap_v4l.cpp index 045c6f889..c9fca0581 100644 --- a/modules/highgui/src/cap_v4l.cpp +++ b/modules/highgui/src/cap_v4l.cpp @@ -325,7 +325,6 @@ typedef struct CvCaptureCAM_V4L struct v4l2_control control; enum v4l2_buf_type type; struct v4l2_queryctrl queryctrl; - struct v4l2_querymenu querymenu; struct timeval timestamp; @@ -641,24 +640,6 @@ static int autosetup_capture_mode_v4l(CvCaptureCAM_V4L* capture) #ifdef HAVE_CAMV4L2 -static void v4l2_scan_controls_enumerate_menu(CvCaptureCAM_V4L* capture) -{ -// printf (" Menu items:\n"); - CLEAR (capture->querymenu); - capture->querymenu.id = capture->queryctrl.id; - for (capture->querymenu.index = capture->queryctrl.minimum; - (int)capture->querymenu.index <= capture->queryctrl.maximum; - capture->querymenu.index++) - { - if (0 == ioctl (capture->deviceHandle, VIDIOC_QUERYMENU, - &capture->querymenu)) - { -// printf (" %s\n", capture->querymenu.name); - } else { - perror ("VIDIOC_QUERYMENU"); - } - } -} static void v4l2_scan_controls(CvCaptureCAM_V4L* capture) { @@ -723,8 +704,6 @@ static void v4l2_scan_controls(CvCaptureCAM_V4L* capture) capture->v4l2_exposure_max = capture->queryctrl.maximum; } - if (capture->queryctrl.type == V4L2_CTRL_TYPE_MENU) - v4l2_scan_controls_enumerate_menu(capture); } else { @@ -793,9 +772,6 @@ static void v4l2_scan_controls(CvCaptureCAM_V4L* capture) capture->v4l2_exposure_max = capture->queryctrl.maximum; } - if (capture->queryctrl.type == V4L2_CTRL_TYPE_MENU) - v4l2_scan_controls_enumerate_menu(capture); - } else { if (errno == EINVAL) From 4aa891e77384e88a4e809f92fa2bc9b50f67a07c Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Wed, 5 Mar 2014 17:57:06 +0400 Subject: [PATCH 07/10] Remove clReleaseDevice call (workaround for pure virtual call on Windows) --- modules/ocl/src/cl_context.cpp | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/modules/ocl/src/cl_context.cpp b/modules/ocl/src/cl_context.cpp index 37d6bb33b..4edf1f7ff 100644 --- a/modules/ocl/src/cl_context.cpp +++ b/modules/ocl/src/cl_context.cpp @@ -596,9 +596,16 @@ protected: CV_Assert(this != currentContext); #ifdef CL_VERSION_1_2 - if (supportsFeature(FEATURE_CL_VER_1_2)) +#ifdef WIN32 + // if process is on termination stage (ExitProcess was called and other threads were terminated) + // then disable command queue release because it may cause program hang + if (!__termination) +#endif { - openCLSafeCall(clReleaseDevice(clDeviceID)); + if (supportsFeature(FEATURE_CL_VER_1_2)) + { + openCLSafeCall(clReleaseDevice(clDeviceID)); + } } #endif if (deviceInfoImpl._id < 0) // not in the global registry, so we should cleanup it From f7a474180be25f472a1e68426f803280454edc56 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Thu, 6 Mar 2014 09:24:02 +0400 Subject: [PATCH 08/10] tuned some tests --- modules/ocl/perf/perf_filters.cpp | 2 +- modules/ocl/perf/perf_kalman.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/modules/ocl/perf/perf_filters.cpp b/modules/ocl/perf/perf_filters.cpp index 70e51b466..b3ffc51b3 100644 --- a/modules/ocl/perf/perf_filters.cpp +++ b/modules/ocl/perf/perf_filters.cpp @@ -298,7 +298,7 @@ OCL_PERF_TEST_P(ScharrFixture, Scharr, oclDst.download(dst); - SANITY_CHECK(dst, 3e-3); + SANITY_CHECK(dst, 1e-2); } else if (RUN_PLAIN_IMPL) { diff --git a/modules/ocl/perf/perf_kalman.cpp b/modules/ocl/perf/perf_kalman.cpp index e384fcd69..946444ad9 100644 --- a/modules/ocl/perf/perf_kalman.cpp +++ b/modules/ocl/perf/perf_kalman.cpp @@ -46,7 +46,7 @@ #include "perf_precomp.hpp" -//#ifdef HAVE_CLAMDBLAS +#ifdef HAVE_CLAMDBLAS using namespace perf; using namespace std; @@ -100,4 +100,4 @@ PERF_TEST_P(KalmanFilterFixture, KalmanFilter, SANITY_CHECK(statePre_); } -//#endif // HAVE_CLAMDBLAS +#endif // HAVE_CLAMDBLAS From 806e9241a6c0255cde117e24e3a9de6fd315443e Mon Sep 17 00:00:00 2001 From: Alexander Karsakov Date: Thu, 6 Mar 2014 11:16:38 +0400 Subject: [PATCH 09/10] Clarifying comments --- modules/ocl/src/cl_context.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/modules/ocl/src/cl_context.cpp b/modules/ocl/src/cl_context.cpp index 4edf1f7ff..156fa9916 100644 --- a/modules/ocl/src/cl_context.cpp +++ b/modules/ocl/src/cl_context.cpp @@ -598,7 +598,7 @@ protected: #ifdef CL_VERSION_1_2 #ifdef WIN32 // if process is on termination stage (ExitProcess was called and other threads were terminated) - // then disable command queue release because it may cause program hang + // then disable device release because it may cause program hang if (!__termination) #endif { @@ -624,7 +624,7 @@ protected: #ifdef WIN32 // if process is on termination stage (ExitProcess was called and other threads were terminated) - // then disable command queue release because it may cause program hang + // then disable context release because it may cause program hang if (!__termination) #endif { From 3eff05e3eb9297d0be9bb72f1131eee5947c209c Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Fri, 7 Mar 2014 14:35:51 +0400 Subject: [PATCH 10/10] added performance tests for cv::ocl::calcHist --- modules/ocl/perf/perf_imgproc.cpp | 39 +++++++++++++++++++++++++++++++ 1 file changed, 39 insertions(+) diff --git a/modules/ocl/perf/perf_imgproc.cpp b/modules/ocl/perf/perf_imgproc.cpp index e6f6a0582..051ff2dba 100644 --- a/modules/ocl/perf/perf_imgproc.cpp +++ b/modules/ocl/perf/perf_imgproc.cpp @@ -81,6 +81,45 @@ OCL_PERF_TEST_P(EqualizeHistFixture, EqualizeHist, OCL_TEST_SIZES) OCL_PERF_ELSE } +///////////// CalcHist //////////////////////// + +typedef TestBaseWithParam CalcHistFixture; + +OCL_PERF_TEST_P(CalcHistFixture, CalcHist, OCL_TEST_SIZES) +{ + const Size srcSize = GetParam(); + const std::vector channels(1, 0); + std::vector ranges(2); + std::vector histSize(1, 256); + ranges[0] = 0; + ranges[1] = 256; + + Mat src(srcSize, CV_8UC1), dst(srcSize, CV_32FC1); + declare.in(src, WARMUP_RNG).out(dst); + + if (RUN_OCL_IMPL) + { + ocl::oclMat oclSrc(src), oclDst(srcSize, CV_32SC1); + + OCL_TEST_CYCLE() cv::ocl::calcHist(oclSrc, oclDst); + + oclDst.download(dst); + SANITY_CHECK(dst); + } + else if (RUN_PLAIN_IMPL) + { + TEST_CYCLE() cv::calcHist(std::vector(1, src), channels, + noArray(), dst, histSize, ranges, false); + + dst.convertTo(dst, CV_32S); + dst = dst.reshape(1, 1); + + SANITY_CHECK(dst); + } + else + OCL_PERF_ELSE +} + /////////// CopyMakeBorder ////////////////////// CV_ENUM(Border, BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101)