From: Andrey Kamaev Date: Wed, 17 Apr 2013 08:07:17 +0000 (+0400) Subject: Merge branch '2.4' X-Git-Tag: submit/tizen_ivi/20141117.190038~2^2~1073 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=8fdab9f631381ad45986c43f48bdb6f533dba541;p=profile%2Fivi%2Fopencv.git Merge branch '2.4' --- 8fdab9f631381ad45986c43f48bdb6f533dba541 diff --cc modules/features2d/perf/perf_fast.cpp index 465bef7,f6a47fe..fe73961 --- a/modules/features2d/perf/perf_fast.cpp +++ b/modules/features2d/perf/perf_fast.cpp @@@ -18,10 -18,10 +18,10 @@@ typedef perf::TestBaseWithParam(GetParam())); + string filename = getDataPath(get<0>(GetParam())); int type = get<1>(GetParam()); Mat frame = imread(filename, IMREAD_GRAYSCALE); diff --cc modules/gpu/perf/perf_core.cpp index 0b20214,3042bea..829637f --- a/modules/gpu/perf/perf_core.cpp +++ b/modules/gpu/perf/perf_core.cpp @@@ -1304,10 -1303,8 +1303,10 @@@ PERF_TEST_P(Sz_3Depth, Core_AddWeighted ////////////////////////////////////////////////////////////////////// // GEMM - CV_FLAGS(GemmFlags, 0, cv::GEMM_1_T, cv::GEMM_2_T, cv::GEMM_3_T) + CV_FLAGS(GemmFlags, 0, GEMM_1_T, GEMM_2_T, GEMM_3_T) -#define ALL_GEMM_FLAGS Values(0, CV_GEMM_A_T, CV_GEMM_B_T, CV_GEMM_C_T, CV_GEMM_A_T | CV_GEMM_B_T, CV_GEMM_A_T | CV_GEMM_C_T, CV_GEMM_A_T | CV_GEMM_B_T | CV_GEMM_C_T) +#define ALL_GEMM_FLAGS Values(0, (int)cv::GEMM_1_T, (int)cv::GEMM_2_T, (int)cv::GEMM_3_T, \ + (int)cv::GEMM_1_T | cv::GEMM_2_T, (int)cv::GEMM_1_T | cv::GEMM_3_T, \ + (int)cv::GEMM_1_T | cv::GEMM_2_T | cv::GEMM_3_T) DEF_PARAM_TEST(Sz_Type_Flags, cv::Size, MatType, GemmFlags); @@@ -2073,12 -2069,9 +2071,9 @@@ PERF_TEST_P(Sz_Depth, Core_CountNonZero ////////////////////////////////////////////////////////////////////// // Reduce - CV_ENUM(ReduceCode, cv::REDUCE_SUM, cv::REDUCE_AVG, cv::REDUCE_MAX, cv::REDUCE_MIN) - #define ALL_REDUCE_CODES ValuesIn(ReduceCode::all()) - enum {Rows = 0, Cols = 1}; -CV_ENUM(ReduceCode, CV_REDUCE_SUM, CV_REDUCE_AVG, CV_REDUCE_MAX, CV_REDUCE_MIN) ++CV_ENUM(ReduceCode, REDUCE_SUM, REDUCE_AVG, REDUCE_MAX, REDUCE_MIN) CV_ENUM(ReduceDim, Rows, Cols) - #define ALL_REDUCE_DIMS ValuesIn(ReduceDim::all()) DEF_PARAM_TEST(Sz_Depth_Cn_Code_Dim, cv::Size, MatDepth, MatCn, ReduceCode, ReduceDim); diff --cc modules/imgproc/perf/perf_cvt_color.cpp index b7eb71f,9b87afe..df1456c --- a/modules/imgproc/perf/perf_cvt_color.cpp +++ b/modules/imgproc/perf/perf_cvt_color.cpp @@@ -338,28 -338,3 +338,28 @@@ PERF_TEST_P(Size_CvtMode3, cvtColorRGB2 SANITY_CHECK(dst, 1); } + +CV_ENUM(EdgeAwareBayerMode, COLOR_BayerBG2BGR_EA, COLOR_BayerGB2BGR_EA, COLOR_BayerRG2BGR_EA, COLOR_BayerGR2BGR_EA) + +typedef std::tr1::tuple EdgeAwareParams; +typedef perf::TestBaseWithParam EdgeAwareDemosaicingTest; + +PERF_TEST_P(EdgeAwareDemosaicingTest, demosaicingEA, + testing::Combine( + testing::Values(szVGA, sz720p, sz1080p, Size(130, 60)), - testing::ValuesIn(EdgeAwareBayerMode::all()) ++ EdgeAwareBayerMode::all() + ) + ) +{ + Size sz = get<0>(GetParam()); + int mode = get<1>(GetParam()); + + Mat src(sz, CV_8UC1); + Mat dst(sz, CV_8UC3); + + declare.in(src, WARMUP_RNG).out(dst); + + TEST_CYCLE() cvtColor(src, dst, mode, 3); + + SANITY_CHECK(dst, 1); +} diff --cc modules/ocl/include/opencv2/ocl/private/util.hpp index 786f2d6,081d234..77d52bf --- a/modules/ocl/include/opencv2/ocl/private/util.hpp +++ b/modules/ocl/include/opencv2/ocl/private/util.hpp @@@ -68,8 -68,7 +68,8 @@@ namespace c void CV_EXPORTS openCLMallocPitch(Context *clCxt, void **dev_ptr, size_t *pitch, size_t widthInBytes, size_t height); void CV_EXPORTS openCLMallocPitchEx(Context *clCxt, void **dev_ptr, size_t *pitch, - size_t widthInBytes, size_t height, - size_t widthInBytes, size_t height, DevMemRW rw_type, DevMemType mem_type); ++ size_t widthInBytes, size_t height, + DevMemRW rw_type, DevMemType mem_type, void* hptr = 0); void CV_EXPORTS openCLMemcpy2D(Context *clCxt, void *dst, size_t dpitch, const void *src, size_t spitch, size_t width, size_t height, openCLMemcpyKind kind, int channels = -1); diff --cc modules/ocl/src/brute_force_matcher.cpp index 74e01b5,9c4a217..7f8b545 --- a/modules/ocl/src/brute_force_matcher.cpp +++ b/modules/ocl/src/brute_force_matcher.cpp @@@ -60,6 -62,17 +61,17 @@@ namespace c } } + static const int OPT_SIZE = 100; + + static const char * T_ARR [] = { - "uchar", - "char", - "ushort", - "short", - "int", - "float -D T_FLOAT", ++ "uchar", ++ "char", ++ "ushort", ++ "short", ++ "int", ++ "float -D T_FLOAT", + "double"}; + template < int BLOCK_SIZE, int MAX_DESC_LEN/*, typename Mask*/ > void matchUnrolledCached(const oclMat &query, const oclMat &train, const oclMat &/*mask*/, const oclMat &trainIdx, const oclMat &distance, int distType) @@@ -71,28 -83,30 +82,30 @@@ const size_t smemSize = (BLOCK_SIZE * (MAX_DESC_LEN >= 2 * BLOCK_SIZE ? MAX_DESC_LEN : 2 * BLOCK_SIZE) + BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); int block_size = BLOCK_SIZE; int m_size = MAX_DESC_LEN; - vector< pair > args; + std::vector< std::pair > args; + char opt [OPT_SIZE] = ""; - sprintf(opt, - "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d -D MAX_DESC_LEN=%d", ++ sprintf(opt, ++ "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d -D MAX_DESC_LEN=%d", + T_ARR[query.depth()], distType, block_size, m_size); + if(globalSize[0] != 0) { - args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&train.data )); - //args.push_back( make_pair( sizeof(cl_mem), (void *)&mask.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&distance.data )); - args.push_back( make_pair( smemSize, (void *)NULL)); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.step )); - - std::string kernelName = "BruteForceMatch_UnrollMatch"; + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&query.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&train.data )); + //args.push_back( std::make_pair( sizeof(cl_mem), (void *)&mask.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&distance.data )); + args.push_back( std::make_pair( smemSize, (void *)NULL)); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&block_size )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&m_size )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.step )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&distType )); + + String kernelName = "BruteForceMatch_UnrollMatch"; - openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth()); + openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt); } } @@@ -112,27 -125,29 +124,29 @@@ void match(const oclMat &query, const o size_t localSize[] = {BLOCK_SIZE, BLOCK_SIZE, 1}; const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); int block_size = BLOCK_SIZE; - vector< pair > args; + std::vector< std::pair > args; + char opt [OPT_SIZE] = ""; - sprintf(opt, - "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d", ++ sprintf(opt, ++ "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d", + T_ARR[query.depth()], distType, block_size); if(globalSize[0] != 0) { - args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&train.data )); - //args.push_back( make_pair( sizeof(cl_mem), (void *)&mask.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&distance.data )); - args.push_back( make_pair( smemSize, (void *)NULL)); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.step )); - - std::string kernelName = "BruteForceMatch_Match"; + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&query.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&train.data )); + //args.push_back( std::make_pair( sizeof(cl_mem), (void *)&mask.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&distance.data )); + args.push_back( std::make_pair( smemSize, (void *)NULL)); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&block_size )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.step )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&distType )); + + String kernelName = "BruteForceMatch_Match"; - openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth()); + openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt); } } @@@ -154,32 -168,34 +167,34 @@@ void matchUnrolledCached(const oclMat & const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); int block_size = BLOCK_SIZE; int m_size = MAX_DESC_LEN; - vector< pair > args; + std::vector< std::pair > args; + char opt [OPT_SIZE] = ""; - sprintf(opt, - "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d -D MAX_DESC_LEN=%d", ++ sprintf(opt, ++ "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d -D MAX_DESC_LEN=%d", + T_ARR[query.depth()], distType, block_size, m_size); + if(globalSize[0] != 0) { - args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&train.data )); - args.push_back( make_pair( sizeof(cl_float), (void *)&maxDistance )); - //args.push_back( make_pair( sizeof(cl_mem), (void *)&mask.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&distance.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&nMatches.data )); - args.push_back( make_pair( smemSize, (void *)NULL)); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&trainIdx.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.step )); - args.push_back( make_pair( sizeof(cl_int), (void *)&trainIdx.step )); - - std::string kernelName = "BruteForceMatch_RadiusUnrollMatch"; + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&query.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&train.data )); + args.push_back( std::make_pair( sizeof(cl_float), (void *)&maxDistance )); + //args.push_back( std::make_pair( sizeof(cl_mem), (void *)&mask.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&distance.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&nMatches.data )); + args.push_back( std::make_pair( smemSize, (void *)NULL)); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&block_size )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&m_size )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&trainIdx.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.step )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&trainIdx.step )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&distType )); + + String kernelName = "BruteForceMatch_RadiusUnrollMatch"; - openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth()); + openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt); } } @@@ -194,31 -209,34 +208,34 @@@ void radius_match(const oclMat &query, size_t localSize[] = {BLOCK_SIZE, BLOCK_SIZE, 1}; const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); int block_size = BLOCK_SIZE; - vector< pair > args; + std::vector< std::pair > args; + char opt [OPT_SIZE] = ""; - sprintf(opt, - "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d", ++ sprintf(opt, ++ "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d", + T_ARR[query.depth()], distType, block_size); + if(globalSize[0] != 0) { - args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&train.data )); - args.push_back( make_pair( sizeof(cl_float), (void *)&maxDistance )); - //args.push_back( make_pair( sizeof(cl_mem), (void *)&mask.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&distance.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&nMatches.data )); - args.push_back( make_pair( smemSize, (void *)NULL)); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&trainIdx.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.step )); - args.push_back( make_pair( sizeof(cl_int), (void *)&trainIdx.step )); - - std::string kernelName = "BruteForceMatch_RadiusMatch"; + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&query.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&train.data )); + args.push_back( std::make_pair( sizeof(cl_float), (void *)&maxDistance )); + //args.push_back( std::make_pair( sizeof(cl_mem), (void *)&mask.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&distance.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&nMatches.data )); + args.push_back( std::make_pair( smemSize, (void *)NULL)); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&block_size )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&trainIdx.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.step )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&trainIdx.step )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&distType )); + + String kernelName = "BruteForceMatch_RadiusMatch"; - openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth()); + openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt); } } @@@ -291,28 -309,30 +308,30 @@@ void knn_matchUnrolledCached(const oclM const size_t smemSize = (BLOCK_SIZE * (MAX_DESC_LEN >= BLOCK_SIZE ? MAX_DESC_LEN : BLOCK_SIZE) + BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); int block_size = BLOCK_SIZE; int m_size = MAX_DESC_LEN; - vector< pair > args; + std::vector< std::pair > args; + char opt [OPT_SIZE] = ""; - sprintf(opt, - "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d -D MAX_DESC_LEN=%d", ++ sprintf(opt, ++ "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d -D MAX_DESC_LEN=%d", + T_ARR[query.depth()], distType, block_size, m_size); + if(globalSize[0] != 0) { - args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&train.data )); - //args.push_back( make_pair( sizeof(cl_mem), (void *)&mask.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&distance.data )); - args.push_back( make_pair( smemSize, (void *)NULL)); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.step )); - - std::string kernelName = "BruteForceMatch_knnUnrollMatch"; + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&query.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&train.data )); + //args.push_back( std::make_pair( sizeof(cl_mem), (void *)&mask.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&distance.data )); + args.push_back( std::make_pair( smemSize, (void *)NULL)); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&block_size )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&m_size )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.step )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&distType )); + + String kernelName = "BruteForceMatch_knnUnrollMatch"; - openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth()); + openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt); } } @@@ -325,27 -345,30 +344,30 @@@ void knn_match(const oclMat &query, con size_t localSize[] = {BLOCK_SIZE, BLOCK_SIZE, 1}; const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); int block_size = BLOCK_SIZE; - vector< pair > args; + std::vector< std::pair > args; + char opt [OPT_SIZE] = ""; - sprintf(opt, - "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d", ++ sprintf(opt, ++ "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d", + T_ARR[query.depth()], distType, block_size); + if(globalSize[0] != 0) { - args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&train.data )); - //args.push_back( make_pair( sizeof(cl_mem), (void *)&mask.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&distance.data )); - args.push_back( make_pair( smemSize, (void *)NULL)); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.step )); - - std::string kernelName = "BruteForceMatch_knnMatch"; + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&query.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&train.data )); + //args.push_back( std::make_pair( sizeof(cl_mem), (void *)&mask.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&distance.data )); + args.push_back( std::make_pair( smemSize, (void *)NULL)); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&block_size )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.step )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&distType )); + + String kernelName = "BruteForceMatch_knnMatch"; - openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth()); + openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt); } } @@@ -358,27 -381,31 +380,31 @@@ void calcDistanceUnrolled(const oclMat const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); int block_size = BLOCK_SIZE; int m_size = MAX_DESC_LEN; - vector< pair > args; + std::vector< std::pair > args; + char opt [OPT_SIZE] = ""; - sprintf(opt, - "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d -D MAX_DESC_LEN=%d", ++ sprintf(opt, ++ "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d -D MAX_DESC_LEN=%d", + T_ARR[query.depth()], distType, block_size, m_size); + if(globalSize[0] != 0) { - args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&train.data )); - //args.push_back( make_pair( sizeof(cl_mem), (void *)&mask.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&allDist.data )); - args.push_back( make_pair( smemSize, (void *)NULL)); - args.push_back( make_pair( sizeof(cl_int), (void *)&block_size )); - args.push_back( make_pair( sizeof(cl_int), (void *)&m_size )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.step )); - - std::string kernelName = "BruteForceMatch_calcDistanceUnrolled"; + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&query.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&train.data )); + //args.push_back( std::make_pair( sizeof(cl_mem), (void *)&mask.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&allDist.data )); + args.push_back( std::make_pair( smemSize, (void *)NULL)); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&block_size )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&m_size )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.step )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&distType )); + + String kernelName = "BruteForceMatch_calcDistanceUnrolled"; - openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth()); + openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt); } } @@@ -390,26 -417,30 +416,30 @@@ void calcDistance(const oclMat &query, size_t localSize[] = {BLOCK_SIZE, BLOCK_SIZE, 1}; const size_t smemSize = (2 * BLOCK_SIZE * BLOCK_SIZE) * sizeof(int); int block_size = BLOCK_SIZE; - vector< pair > args; + std::vector< std::pair > args; + char opt [OPT_SIZE] = ""; - sprintf(opt, - "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d", ++ sprintf(opt, ++ "-D T=%s -D DIST_TYPE=%d -D BLOCK_SIZE=%d", + T_ARR[query.depth()], distType, block_size); + if(globalSize[0] != 0) { - args.push_back( make_pair( sizeof(cl_mem), (void *)&query.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&train.data )); - //args.push_back( make_pair( sizeof(cl_mem), (void *)&mask.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&allDist.data )); - args.push_back( make_pair( smemSize, (void *)NULL)); - args.push_back( make_pair( sizeof(cl_int), (void *)&block_size )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.rows )); - args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols )); - args.push_back( make_pair( sizeof(cl_int), (void *)&query.step )); - - std::string kernelName = "BruteForceMatch_calcDistance"; + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&query.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&train.data )); + //args.push_back( std::make_pair( sizeof(cl_mem), (void *)&mask.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&allDist.data )); + args.push_back( std::make_pair( smemSize, (void *)NULL)); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&block_size )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.rows )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.cols )); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.step )); - args.push_back( std::make_pair( sizeof(cl_int), (void *)&distType )); + + String kernelName = "BruteForceMatch_calcDistance"; - openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, query.depth()); + openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1, opt); } } @@@ -460,18 -491,18 +490,18 @@@ void findKnnMatch(int k, const oclMat & for (int i = 0; i < k; ++i) { - vector< pair > args; + std::vector< std::pair > args; - args.push_back( make_pair( sizeof(cl_mem), (void *)&allDist.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&distance.data )); - args.push_back( make_pair( sizeof(cl_mem), (void *)&i)); - args.push_back( make_pair( sizeof(cl_int), (void *)&block_size )); - //args.push_back( make_pair( sizeof(cl_int), (void *)&train.rows )); - //args.push_back( make_pair( sizeof(cl_int), (void *)&train.cols )); - //args.push_back( make_pair( sizeof(cl_int), (void *)&query.step )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&allDist.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&trainIdx.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&distance.data )); + args.push_back( std::make_pair( sizeof(cl_mem), (void *)&i)); + args.push_back( std::make_pair( sizeof(cl_int), (void *)&block_size )); + //args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.rows )); + //args.push_back( std::make_pair( sizeof(cl_int), (void *)&train.cols )); + //args.push_back( std::make_pair( sizeof(cl_int), (void *)&query.step )); - openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, trainIdx.depth(), -1); + openCLExecuteKernel(ctx, &brute_force_match, kernelName, globalSize, localSize, args, -1, -1); } } @@@ -530,28 -561,19 +560,19 @@@ void cv::ocl::BruteForceMatcher_OCL_bas { if (query.empty() || train.empty()) return; - + - // match1 doesn't support signed char type, match2 only support float, hamming support uchar, ushort and int - int callType = query.depth(); - if (callType != 5) - CV_Error(Error::StsUnsupportedFormat, "BruteForceMatch OpenCL only support float type query!\n"); - - if ((distType == 0 && callType == 1 ) || (distType == 1 && callType != 5) || (distType == 2 && (callType != 0 - || callType != 2 || callType != 4))) - { - CV_Error(Error::BadDepth, "BruteForceMatch OpenCL only support float type query!\n"); - } - CV_Assert(query.channels() == 1 && query.depth() < CV_64F); CV_Assert(train.cols == query.cols && train.type() == query.type()); - trainIdx.create(1, query.rows, CV_32S); - distance.create(1, query.rows, CV_32F); + ensureSizeIsEnough(1, query.rows, CV_32S, trainIdx); + ensureSizeIsEnough(1, query.rows, CV_32F, distance); matchDispatcher(query, train, mask, trainIdx, distance, distType); + + return; } -void cv::ocl::BruteForceMatcher_OCL_base::matchDownload(const oclMat &trainIdx, const oclMat &distance, vector &matches) +void cv::ocl::BruteForceMatcher_OCL_base::matchDownload(const oclMat &trainIdx, const oclMat &distance, std::vector &matches) { if (trainIdx.empty() || distance.empty()) return; @@@ -592,9 -614,9 +613,9 @@@ void cv::ocl::BruteForceMatcher_OCL_bas } } -void cv::ocl::BruteForceMatcher_OCL_base::match(const oclMat &query, const oclMat &train, vector &matches, const oclMat &mask) +void cv::ocl::BruteForceMatcher_OCL_base::match(const oclMat &query, const oclMat &train, std::vector &matches, const oclMat &mask) { - CV_Assert(mask.empty()); // mask is not supported at the moment + assert(mask.empty()); // mask is not supported at the moment oclMat trainIdx, distance; matchSingle(query, train, trainIdx, distance, mask); matchDownload(trainIdx, distance, matches); @@@ -650,27 -672,20 +671,20 @@@ void cv::ocl::BruteForceMatcher_OCL_bas if (query.empty() || trainCollection.empty()) return; - // match1 doesn't support signed char type, match2 only support float, hamming support uchar, ushort and int - int callType = query.depth(); - if (callType != 5) - CV_Error(Error::StsUnsupportedFormat, "BruteForceMatch OpenCL only support float type query!\n"); - - if ((distType == 0 && callType == 1 ) || (distType == 1 && callType != 5) || (distType == 2 && (callType != 0 - || callType != 2 || callType != 4))) - { - CV_Error(Error::BadDepth, "BruteForceMatch OpenCL only support float type query!\n"); - } - CV_Assert(query.channels() == 1 && query.depth() < CV_64F); - + - trainIdx.create(1, query.rows, CV_32S); - imgIdx.create(1, query.rows, CV_32S); - distance.create(1, query.rows, CV_32F); + const int nQuery = query.rows; + + ensureSizeIsEnough(1, nQuery, CV_32S, trainIdx); + ensureSizeIsEnough(1, nQuery, CV_32S, imgIdx); + ensureSizeIsEnough(1, nQuery, CV_32F, distance); matchDispatcher(query, (const oclMat *)trainCollection.ptr(), trainCollection.cols, masks, trainIdx, imgIdx, distance, distType); + + return; } -void cv::ocl::BruteForceMatcher_OCL_base::matchDownload(const oclMat &trainIdx, const oclMat &imgIdx, const oclMat &distance, vector &matches) +void cv::ocl::BruteForceMatcher_OCL_base::matchDownload(const oclMat &trainIdx, const oclMat &imgIdx, const oclMat &distance, std::vector &matches) { if (trainIdx.empty() || imgIdx.empty() || distance.empty()) return; @@@ -766,9 -772,11 +771,11 @@@ void cv::ocl::BruteForceMatcher_OCL_bas trainIdx.setTo(Scalar::all(-1)); kmatchDispatcher(query, train, k, mask, trainIdx, distance, allDist, distType); + + return; } -void cv::ocl::BruteForceMatcher_OCL_base::knnMatchDownload(const oclMat &trainIdx, const oclMat &distance, vector< vector > &matches, bool compactResult) +void cv::ocl::BruteForceMatcher_OCL_base::knnMatchDownload(const oclMat &trainIdx, const oclMat &distance, std::vector< std::vector > &matches, bool compactResult) { if (trainIdx.empty() || distance.empty()) return; @@@ -972,7 -961,7 +960,7 @@@ void cv::ocl::BruteForceMatcher_OCL_bas temp.reserve(2 * k); matches.resize(query.rows); - std::for_each(matches.begin(), matches.end(), std::bind2nd(std::mem_fun_ref(&std::vector::reserve), k)); - for_each(matches.begin(), matches.end(), bind2nd(mem_fun_ref(&vector::reserve), k)); ++ for_each(matches.begin(), matches.end(), bind2nd(mem_fun_ref(&std::vector::reserve), k)); for (size_t imgIdx = 0, size = trainDescCollection.size(); imgIdx < size; ++imgIdx) { @@@ -996,7 -985,7 +984,7 @@@ if (compactResult) { - std::vector< std::vector >::iterator new_end = std::remove_if(matches.begin(), matches.end(), std::mem_fun_ref(&std::vector::empty)); - vector< vector >::iterator new_end = remove_if(matches.begin(), matches.end(), mem_fun_ref(&vector::empty)); ++ std::vector< std::vector >::iterator new_end = remove_if(matches.begin(), matches.end(), mem_fun_ref(&std::vector::empty)); matches.erase(new_end, matches.end()); } } diff --cc modules/ocl/src/opencl/brute_force_match.cl index fc7e515,8dcb9d2..a05c98e --- a/modules/ocl/src/opencl/brute_force_match.cl +++ b/modules/ocl/src/opencl/brute_force_match.cl @@@ -405,7 -425,8 +425,8 @@@ __kernel void BruteForceMatch_RadiusMat barrier(CLK_LOCAL_MEM_FENCE); } - if (queryIdx < query_rows && trainIdx < train_rows && result < maxDistance/* && mask(queryIdx, trainIdx)*/) - if (queryIdx < query_rows && trainIdx < train_rows && ++ if (queryIdx < query_rows && trainIdx < train_rows && + convert_float(result) < maxDistance/* && mask(queryIdx, trainIdx)*/) { unsigned int ind = atom_inc(nMatches + queryIdx); diff --cc modules/ocl/test/test_brute_force_matcher.cpp index 71b299e,793cbd4..1fbbd99 --- a/modules/ocl/test/test_brute_force_matcher.cpp +++ b/modules/ocl/test/test_brute_force_matcher.cpp @@@ -189,10 -185,22 +185,22 @@@ namespac ASSERT_EQ(0, badCount); } - INSTANTIATE_TEST_CASE_P(OCL_Features2D, BruteForceMatcher, + INSTANTIATE_TEST_CASE_P(OCL_Features2D, BruteForceMatcher, testing::Combine( - testing::Values(DistType(cv::ocl::BruteForceMatcher_OCL_base::L1Dist), DistType(cv::ocl::BruteForceMatcher_OCL_base::L2Dist)), - testing::Values(DescriptorSize(57), DescriptorSize(64), DescriptorSize(83), DescriptorSize(128), DescriptorSize(179), DescriptorSize(256), DescriptorSize(304)))); - + testing::Values( + DistType(cv::ocl::BruteForceMatcher_OCL_base::L1Dist), - DistType(cv::ocl::BruteForceMatcher_OCL_base::L2Dist)/*, ++ DistType(cv::ocl::BruteForceMatcher_OCL_base::L2Dist)/*, + DistType(cv::ocl::BruteForceMatcher_OCL_base::HammingDist)*/ + ), + testing::Values( - DescriptorSize(57), - DescriptorSize(64), - DescriptorSize(83), - DescriptorSize(128), - DescriptorSize(179), - DescriptorSize(256), ++ DescriptorSize(57), ++ DescriptorSize(64), ++ DescriptorSize(83), ++ DescriptorSize(128), ++ DescriptorSize(179), ++ DescriptorSize(256), + DescriptorSize(304)) + ) + ); } // namespace #endif diff --cc modules/ocl/test/utility.hpp index 838f65a,9e8a087..947cb37 --- a/modules/ocl/test/utility.hpp +++ b/modules/ocl/test/utility.hpp @@@ -135,27 -133,18 +135,18 @@@ void PrintTo(const Inverse &useRoi, std enum {FLIP_BOTH = 0, FLIP_X = 1, FLIP_Y = -1}; CV_ENUM(FlipCode, FLIP_BOTH, FLIP_X, FLIP_Y) - CV_ENUM(ReduceOp, cv::REDUCE_SUM, cv::REDUCE_AVG, cv::REDUCE_MAX, cv::REDUCE_MIN) - - CV_FLAGS(GemmFlags, cv::GEMM_1_T, cv::GEMM_2_T, cv::GEMM_3_T); - - CV_ENUM(MorphOp, cv::MORPH_OPEN, cv::MORPH_CLOSE, cv::MORPH_GRADIENT, cv::MORPH_TOPHAT, cv::MORPH_BLACKHAT) - - CV_ENUM(ThreshOp, cv::THRESH_BINARY, cv::THRESH_BINARY_INV, cv::THRESH_TRUNC, cv::THRESH_TOZERO, cv::THRESH_TOZERO_INV) - - CV_ENUM(Interpolation, cv::INTER_NEAREST, cv::INTER_LINEAR, cv::INTER_CUBIC) - - CV_ENUM(Border, cv::BORDER_REFLECT101, cv::BORDER_REPLICATE, cv::BORDER_CONSTANT, cv::BORDER_REFLECT, cv::BORDER_WRAP) - - CV_FLAGS(WarpFlags, cv::INTER_NEAREST, cv::INTER_LINEAR, cv::INTER_CUBIC, cv::WARP_INVERSE_MAP) - - CV_ENUM(TemplateMethod, cv::TM_SQDIFF, cv::TM_SQDIFF_NORMED, cv::TM_CCORR, cv::TM_CCORR_NORMED, cv::TM_CCOEFF, cv::TM_CCOEFF_NORMED) - - CV_FLAGS(DftFlags, cv::DFT_INVERSE, cv::DFT_SCALE, cv::DFT_ROWS, cv::DFT_COMPLEX_OUTPUT, cv::DFT_REAL_OUTPUT) + CV_ENUM(CmpCode, CMP_EQ, CMP_GT, CMP_GE, CMP_LT, CMP_LE, CMP_NE) + CV_ENUM(NormCode, NORM_INF, NORM_L1, NORM_L2, NORM_TYPE_MASK, NORM_RELATIVE, NORM_MINMAX) -CV_ENUM(ReduceOp, CV_REDUCE_SUM, CV_REDUCE_AVG, CV_REDUCE_MAX, CV_REDUCE_MIN) ++CV_ENUM(ReduceOp, REDUCE_SUM, REDUCE_AVG, REDUCE_MAX, REDUCE_MIN) + CV_ENUM(MorphOp, MORPH_OPEN, MORPH_CLOSE, MORPH_GRADIENT, MORPH_TOPHAT, MORPH_BLACKHAT) + CV_ENUM(ThreshOp, THRESH_BINARY, THRESH_BINARY_INV, THRESH_TRUNC, THRESH_TOZERO, THRESH_TOZERO_INV) + CV_ENUM(Interpolation, INTER_NEAREST, INTER_LINEAR, INTER_CUBIC) + CV_ENUM(Border, BORDER_REFLECT101, BORDER_REPLICATE, BORDER_CONSTANT, BORDER_REFLECT, BORDER_WRAP) + CV_ENUM(TemplateMethod, TM_SQDIFF, TM_SQDIFF_NORMED, TM_CCORR, TM_CCORR_NORMED, TM_CCOEFF, TM_CCOEFF_NORMED) + + CV_FLAGS(GemmFlags, GEMM_1_T, GEMM_2_T, GEMM_3_T); + CV_FLAGS(WarpFlags, INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, WARP_INVERSE_MAP) + CV_FLAGS(DftFlags, DFT_INVERSE, DFT_SCALE, DFT_ROWS, DFT_COMPLEX_OUTPUT, DFT_REAL_OUTPUT) void run_perf_test();