diff --git a/modules/ocl/src/matrix_operations.cpp b/modules/ocl/src/matrix_operations.cpp index 82189b71e5..cd09f44027 100644 --- a/modules/ocl/src/matrix_operations.cpp +++ b/modules/ocl/src/matrix_operations.cpp @@ -382,40 +382,50 @@ void cv::ocl::oclMat::copyTo( oclMat &mat, const oclMat &mask) const /////////////////////////////////////////////////////////////////////////// static void convert_run(const oclMat &src, oclMat &dst, double alpha, double beta) { - string kernelName = "convert_to_S"; - stringstream idxStr; - idxStr << src.depth(); - kernelName += idxStr.str(); + string kernelName = "convert_to"; float alpha_f = alpha, beta_f = beta; + int sdepth = src.depth(), ddepth = dst.depth(); + int sstep1 = (int)src.step1(), dstep1 = (int)dst.step1(); + int cols1 = src.cols * src.oclchannels(); + + char buildOptions[150], convertString[50]; + const char * typeMap[] = { "uchar", "char", "ushort", "short", "int", "float", "double" }; + sprintf(convertString, "convert_%s_sat_rte", typeMap[ddepth]); + sprintf(buildOptions, "-D srcT=%s -D dstT=%s -D convertToDstType=%s", typeMap[sdepth], + typeMap[ddepth], CV_32F == ddepth || ddepth == CV_64F ? "" : convertString); + CV_DbgAssert(src.rows == dst.rows && src.cols == dst.cols); vector > args; - size_t localThreads[3] = {16, 16, 1}; - size_t globalThreads[3]; - globalThreads[0] = (dst.cols + localThreads[0] - 1) / localThreads[0] * localThreads[0]; - globalThreads[1] = (dst.rows + localThreads[1] - 1) / localThreads[1] * localThreads[1]; - globalThreads[2] = 1; - int dststep_in_pixel = dst.step / dst.elemSize(), dstoffset_in_pixel = dst.offset / dst.elemSize(); - int srcstep_in_pixel = src.step / src.elemSize(), srcoffset_in_pixel = src.offset / src.elemSize(); - if(dst.type() == CV_8UC1) - { - globalThreads[0] = ((dst.cols + 4) / 4 + localThreads[0]) / localThreads[0] * localThreads[0]; - } + + size_t localThreads[3] = { 16, 16, 1 }; + size_t globalThreads[3] = { divUp(cols1, localThreads[0]) * localThreads[0], + divUp(dst.rows, localThreads[1]) * localThreads[1], 1 }; + + int doffset1 = dst.offset / dst.elemSize1(); + int soffset1 = src.offset / src.elemSize1(); + args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data )); args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&src.cols )); + args.push_back( make_pair( sizeof(cl_int) , (void *)&cols1 )); args.push_back( make_pair( sizeof(cl_int) , (void *)&src.rows )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&srcstep_in_pixel )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&srcoffset_in_pixel )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dststep_in_pixel )); - args.push_back( make_pair( sizeof(cl_int) , (void *)&dstoffset_in_pixel )); + args.push_back( make_pair( sizeof(cl_int) , (void *)&sstep1 )); + args.push_back( make_pair( sizeof(cl_int) , (void *)&soffset1 )); + args.push_back( make_pair( sizeof(cl_int) , (void *)&dstep1 )); + args.push_back( make_pair( sizeof(cl_int) , (void *)&doffset1 )); args.push_back( make_pair( sizeof(cl_float) , (void *)&alpha_f )); args.push_back( make_pair( sizeof(cl_float) , (void *)&beta_f )); + openCLExecuteKernel(dst.clCxt , &operator_convertTo, kernelName, globalThreads, - localThreads, args, dst.oclchannels(), dst.depth()); + localThreads, args, -1, -1, buildOptions); } void cv::ocl::oclMat::convertTo( oclMat &dst, int rtype, double alpha, double beta ) const { - //cout << "cv::ocl::oclMat::convertTo()" << endl; + if (!clCxt->supportsFeature(Context::CL_DOUBLE) && + (depth() == CV_64F || dst.depth() == CV_64F)) + { + CV_Error(CV_GpuNotSupported, "Selected device don't support double\r\n"); + return; + } bool noScale = fabs(alpha - 1) < std::numeric_limits::epsilon() && fabs(beta) < std::numeric_limits::epsilon(); @@ -425,7 +435,6 @@ void cv::ocl::oclMat::convertTo( oclMat &dst, int rtype, double alpha, double be else rtype = CV_MAKETYPE(CV_MAT_DEPTH(rtype), channels()); - //int scn = channels(); int sdepth = depth(), ddepth = CV_MAT_DEPTH(rtype); if( sdepth == ddepth && noScale ) { @@ -447,7 +456,6 @@ void cv::ocl::oclMat::convertTo( oclMat &dst, int rtype, double alpha, double be /////////////////////////////////////////////////////////////////////////// oclMat &cv::ocl::oclMat::operator = (const Scalar &s) { - //cout << "cv::ocl::oclMat::=" << endl; setTo(s); return *this; } diff --git a/modules/ocl/src/opencl/operator_convertTo.cl b/modules/ocl/src/opencl/operator_convertTo.cl index 1a8dd04b93..278d41f7ca 100644 --- a/modules/ocl/src/opencl/operator_convertTo.cl +++ b/modules/ocl/src/opencl/operator_convertTo.cl @@ -33,352 +33,28 @@ // the use of this software, even if advised of the possibility of such damage. // // -#define F float -#define F2 float2 -#define F4 float4 -__kernel void convert_to_S4_C1_D0( - __global const int* restrict srcMat, - __global uchar* dstMat, - int cols, - int rows, - int srcStep_in_pixel, - int srcoffset_in_pixel, - int dstStep_in_pixel, - int dstoffset_in_pixel, - F alpha, - F beta) -{ - int x=get_global_id(0)<<2; - int y=get_global_id(1); - //int src_addr_start = mad24(y,srcStep_in_pixel,srcoffset_in_pixel); - //int src_addr_end = mad24(y,srcStep_in_pixel,cols+srcoffset_in_pixel); - int off_src = (dstoffset_in_pixel & 3); - int srcidx = mad24(y,srcStep_in_pixel,x+ srcoffset_in_pixel - off_src); - int dst_addr_start = mad24(y,dstStep_in_pixel,dstoffset_in_pixel); - int dst_addr_end = mad24(y,dstStep_in_pixel,cols+dstoffset_in_pixel); - int dstidx = mad24(y,dstStep_in_pixel,x+ dstoffset_in_pixel & (int)0xfffffffc); - if(x+3get_rng(); - cv::Size size(MWIDTH, MHEIGHT); - mat = randomMat(rng, size, type, 5, 16, false); - dst = randomMat(rng, size, type, 5, 16, false); + mat = randomMat(rng, randomSize(MIN_VALUE, MAX_VALUE), src_type, 5, 136, false); + dst = randomMat(rng, use_roi ? randomSize(MIN_VALUE, MAX_VALUE) : mat.size(), dst_type, 5, 136, false); } void random_roi() { -#ifdef RANDOMROI - //randomize ROI - cv::RNG &rng = TS::ptr()->get_rng(); - roicols = rng.uniform(1, mat.cols); - roirows = rng.uniform(1, mat.rows); - srcx = rng.uniform(0, mat.cols - roicols); - srcy = rng.uniform(0, mat.rows - roirows); - dstx = rng.uniform(0, dst.cols - roicols); - dsty = rng.uniform(0, dst.rows - roirows); -#else - roicols = mat.cols; - roirows = mat.rows; - srcx = 0; - srcy = 0; - dstx = 0; - dsty = 0; -#endif + if (use_roi) + { + // randomize ROI + cv::RNG &rng = TS::ptr()->get_rng(); + roicols = rng.uniform(1, MIN_VALUE); + roirows = rng.uniform(1, MIN_VALUE); + srcx = rng.uniform(0, mat.cols - roicols); + srcy = rng.uniform(0, mat.rows - roirows); + dstx = rng.uniform(0, dst.cols - roicols); + dsty = rng.uniform(0, dst.rows - roirows); + } + else + { + roicols = mat.cols; + roirows = mat.rows; + srcx = srcy = 0; + dstx = dsty = 0; + } mat_roi = mat(Rect(srcx, srcy, roicols, roirows)); - dst_roi = dst(Rect(dstx, dsty, roicols, roirows)); + dst_roi = dst(Rect(dstx, dsty, roicols, roirows)); gdst_whole = dst; gdst = gdst_whole(Rect(dstx, dsty, roicols, roirows)); @@ -123,30 +127,28 @@ PARAM_TEST_CASE(ConvertToTestBase, MatType, MatType) } }; - -struct ConvertTo : ConvertToTestBase {}; +typedef ConvertToTestBase ConvertTo; TEST_P(ConvertTo, Accuracy) { - for(int j = 0; j < LOOP_TIMES; j++) + for (int j = 0; j < LOOP_TIMES; j++) { random_roi(); mat_roi.convertTo(dst_roi, dst_type); gmat.convertTo(gdst, dst_type); - EXPECT_MAT_NEAR(dst, Mat(gdst_whole), 0.0); + EXPECT_MAT_NEAR(dst, Mat(gdst_whole), src_depth == CV_64F ? 1.0 : 0.0); + EXPECT_MAT_NEAR(dst_roi, Mat(gdst), src_depth == CV_64F ? 1.0 : 0.0); } } - - - ///////////////////////////////////////////copyto///////////////////////////////////////////////////////////// PARAM_TEST_CASE(CopyToTestBase, MatType, bool) { int type; + bool use_roi; cv::Mat mat; cv::Mat mask; @@ -162,15 +164,15 @@ PARAM_TEST_CASE(CopyToTestBase, MatType, bool) int maskx; int masky; - //src mat with roi + // src mat with roi cv::Mat mat_roi; cv::Mat mask_roi; cv::Mat dst_roi; - //ocl dst mat for testing + // ocl dst mat for testing cv::ocl::oclMat gdst_whole; - //ocl mat with roi + // ocl mat with roi cv::ocl::oclMat gmat; cv::ocl::oclMat gdst; cv::ocl::oclMat gmask; @@ -178,45 +180,45 @@ PARAM_TEST_CASE(CopyToTestBase, MatType, bool) virtual void SetUp() { type = GET_PARAM(0); + use_roi = GET_PARAM(1); cv::RNG &rng = TS::ptr()->get_rng(); cv::Size size(MWIDTH, MHEIGHT); mat = randomMat(rng, size, type, 5, 16, false); - dst = randomMat(rng, size, type, 5, 16, false); + dst = randomMat(rng, size, type, 5, 16, false); mask = randomMat(rng, size, CV_8UC1, 0, 2, false); cv::threshold(mask, mask, 0.5, 255., CV_8UC1); - } void random_roi() { -#ifdef RANDOMROI - //randomize ROI - cv::RNG &rng = TS::ptr()->get_rng(); - roicols = rng.uniform(1, mat.cols); - roirows = rng.uniform(1, mat.rows); - srcx = rng.uniform(0, mat.cols - roicols); - srcy = rng.uniform(0, mat.rows - roirows); - dstx = rng.uniform(0, dst.cols - roicols); - dsty = rng.uniform(0, dst.rows - roirows); - maskx = rng.uniform(0, mask.cols - roicols); - masky = rng.uniform(0, mask.rows - roirows); -#else - roicols = mat.cols; - roirows = mat.rows; - srcx = 0; - srcy = 0; - dstx = 0; - dsty = 0; - maskx = 0; - masky = 0; -#endif + if (use_roi) + { + // randomize ROI + cv::RNG &rng = TS::ptr()->get_rng(); + roicols = rng.uniform(1, mat.cols); + roirows = rng.uniform(1, mat.rows); + srcx = rng.uniform(0, mat.cols - roicols); + srcy = rng.uniform(0, mat.rows - roirows); + dstx = rng.uniform(0, dst.cols - roicols); + dsty = rng.uniform(0, dst.rows - roirows); + maskx = rng.uniform(0, mask.cols - roicols); + masky = rng.uniform(0, mask.rows - roirows); + } + else + { + roicols = mat.cols; + roirows = mat.rows; + srcx = srcy = 0; + dstx = dsty = 0; + maskx = masky = 0; + } mat_roi = mat(Rect(srcx, srcy, roicols, roirows)); mask_roi = mask(Rect(maskx, masky, roicols, roirows)); - dst_roi = dst(Rect(dstx, dsty, roicols, roirows)); + dst_roi = dst(Rect(dstx, dsty, roicols, roirows)); gdst_whole = dst; gdst = gdst_whole(Rect(dstx, dsty, roicols, roirows)); @@ -226,11 +228,11 @@ PARAM_TEST_CASE(CopyToTestBase, MatType, bool) } }; -struct CopyTo : CopyToTestBase {}; +typedef CopyToTestBase CopyTo; TEST_P(CopyTo, Without_mask) { - for(int j = 0; j < LOOP_TIMES; j++) + for (int j = 0; j < LOOP_TIMES; j++) { random_roi(); @@ -243,7 +245,7 @@ TEST_P(CopyTo, Without_mask) TEST_P(CopyTo, With_mask) { - for(int j = 0; j < LOOP_TIMES; j++) + for (int j = 0; j < LOOP_TIMES; j++) { random_roi(); @@ -254,14 +256,13 @@ TEST_P(CopyTo, With_mask) } } - - - -///////////////////////////////////////////copyto///////////////////////////////////////////////////////////// +/////////////////////////////////////////// setTo ///////////////////////////////////////////////////////////// PARAM_TEST_CASE(SetToTestBase, MatType, bool) { int type; + bool use_roi; + cv::Scalar val; cv::Mat mat; @@ -275,20 +276,21 @@ PARAM_TEST_CASE(SetToTestBase, MatType, bool) int maskx; int masky; - //src mat with roi + // src mat with roi cv::Mat mat_roi; cv::Mat mask_roi; - //ocl dst mat for testing + // ocl dst mat for testing cv::ocl::oclMat gmat_whole; - //ocl mat with roi + // ocl mat with roi cv::ocl::oclMat gmat; cv::ocl::oclMat gmask; virtual void SetUp() { type = GET_PARAM(0); + use_roi = GET_PARAM(1); cv::RNG &rng = TS::ptr()->get_rng(); cv::Size size(MWIDTH, MHEIGHT); @@ -298,28 +300,28 @@ PARAM_TEST_CASE(SetToTestBase, MatType, bool) cv::threshold(mask, mask, 0.5, 255., CV_8UC1); val = cv::Scalar(rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0), rng.uniform(-10.0, 10.0)); - } void random_roi() { -#ifdef RANDOMROI - //randomize ROI - cv::RNG &rng = TS::ptr()->get_rng(); - roicols = rng.uniform(1, mat.cols); - roirows = rng.uniform(1, mat.rows); - srcx = rng.uniform(0, mat.cols - roicols); - srcy = rng.uniform(0, mat.rows - roirows); - maskx = rng.uniform(0, mask.cols - roicols); - masky = rng.uniform(0, mask.rows - roirows); -#else - roicols = mat.cols; - roirows = mat.rows; - srcx = 0; - srcy = 0; - maskx = 0; - masky = 0; -#endif + if (use_roi) + { + // randomize ROI + cv::RNG &rng = TS::ptr()->get_rng(); + roicols = rng.uniform(1, mat.cols); + roirows = rng.uniform(1, mat.rows); + srcx = rng.uniform(0, mat.cols - roicols); + srcy = rng.uniform(0, mat.rows - roirows); + maskx = rng.uniform(0, mask.cols - roicols); + masky = rng.uniform(0, mask.rows - roirows); + } + else + { + roicols = mat.cols; + roirows = mat.rows; + srcx = srcy = 0; + maskx = masky = 0; + } mat_roi = mat(Rect(srcx, srcy, roicols, roirows)); mask_roi = mask(Rect(maskx, masky, roicols, roirows)); @@ -331,11 +333,11 @@ PARAM_TEST_CASE(SetToTestBase, MatType, bool) } }; -struct SetTo : SetToTestBase {}; +typedef SetToTestBase SetTo; TEST_P(SetTo, Without_mask) { - for(int j = 0; j < LOOP_TIMES; j++) + for (int j = 0; j < LOOP_TIMES; j++) { random_roi(); @@ -348,7 +350,7 @@ TEST_P(SetTo, Without_mask) TEST_P(SetTo, With_mask) { - for(int j = 0; j < LOOP_TIMES; j++) + for (int j = 0; j < LOOP_TIMES; j++) { random_roi(); @@ -359,105 +361,84 @@ TEST_P(SetTo, With_mask) } } -//convertC3C4 -PARAM_TEST_CASE(convertC3C4, MatType, cv::Size) +// convertC3C4 + +PARAM_TEST_CASE(convertC3C4, MatType, bool) { - int type; - cv::Size ksize; + int depth; + bool use_roi; //src mat - cv::Mat mat1; - cv::Mat dst; + cv::Mat src; // set up roi - int roicols; - int roirows; - int src1x; - int src1y; - int dstx; - int dsty; + int roicols, roirows; + int srcx, srcy; //src mat with roi - cv::Mat mat1_roi; - cv::Mat dst_roi; - - //ocl dst mat for testing - cv::ocl::oclMat gdst_whole; + cv::Mat src_roi; //ocl mat with roi - cv::ocl::oclMat gmat1; - cv::ocl::oclMat gdst; + cv::ocl::oclMat gsrc_roi; virtual void SetUp() { - type = GET_PARAM(0); - ksize = GET_PARAM(1); + depth = GET_PARAM(0); + use_roi = GET_PARAM(1); + int type = CV_MAKE_TYPE(depth, 3); + cv::RNG &rng = TS::ptr()->get_rng(); + src = randomMat(rng, randomSize(MIN_VALUE, MAX_VALUE), type, 0, 40, false); } void random_roi() { -#ifdef RANDOMROI - //randomize ROI - cv::RNG &rng = TS::ptr()->get_rng(); - roicols = rng.uniform(2, mat1.cols); - roirows = rng.uniform(2, mat1.rows); - src1x = rng.uniform(0, mat1.cols - roicols); - src1y = rng.uniform(0, mat1.rows - roirows); - dstx = rng.uniform(0, dst.cols - roicols); - dsty = rng.uniform(0, dst.rows - roirows); -#else - roicols = mat1.cols; - roirows = mat1.rows; - src1x = 0; - src1y = 0; - dstx = 0; - dsty = 0; -#endif - - mat1_roi = mat1(Rect(src1x, src1y, roicols, roirows)); - dst_roi = dst(Rect(dstx, dsty, roicols, roirows)); - - gdst_whole = dst; - gdst = gdst_whole(Rect(dstx, dsty, roicols, roirows)); - - - gmat1 = mat1_roi; + if (use_roi) + { + //randomize ROI + cv::RNG &rng = TS::ptr()->get_rng(); + roicols = rng.uniform(1, src.cols); + roirows = rng.uniform(1, src.rows); + srcx = rng.uniform(0, src.cols - roicols); + srcy = rng.uniform(0, src.rows - roirows); + } + else + { + roicols = src.cols; + roirows = src.rows; + srcx = srcy = 0; + } + + src_roi = src(Rect(srcx, srcy, roicols, roirows)); } - }; TEST_P(convertC3C4, Accuracy) { - cv::RNG &rng = TS::ptr()->get_rng(); - for(int j = 0; j < LOOP_TIMES; j++) + for (int j = 0; j < LOOP_TIMES; j++) { - //random_roi(); - int width = rng.uniform(2, MWIDTH); - int height = rng.uniform(2, MHEIGHT); - cv::Size size(width, height); + random_roi(); - mat1 = randomMat(rng, size, type, 0, 40, false); - gmat1 = mat1; + gsrc_roi = src_roi; - EXPECT_MAT_NEAR(mat1, Mat(gmat1), 0.0); + EXPECT_MAT_NEAR(src_roi, Mat(gsrc_roi), 0.0); } - } INSTANTIATE_TEST_CASE_P(MatrixOperation, ConvertTo, Combine( - Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32SC1, CV_32SC4, CV_32FC1, CV_32FC4), - Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32SC1, CV_32SC4, CV_32FC1, CV_32FC4))); + Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), + Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), + Range(1, 5), Bool())); INSTANTIATE_TEST_CASE_P(MatrixOperation, CopyTo, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32SC1, CV_32SC3, CV_32SC4, CV_32FC1, CV_32FC3, CV_32FC4), - Values(false))); // Values(false) is the reserved parameter + Bool())); INSTANTIATE_TEST_CASE_P(MatrixOperation, SetTo, Combine( Values(CV_8UC1, CV_8UC3, CV_8UC4, CV_32SC1, CV_32SC3, CV_32SC4, CV_32FC1, CV_32FC3, CV_32FC4), - Values(false))); // Values(false) is the reserved parameter + Bool())); INSTANTIATE_TEST_CASE_P(MatrixOperation, convertC3C4, Combine( - Values(CV_8UC3, CV_32SC3, CV_32FC3), - Values(cv::Size()))); + Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), + Bool())); #endif diff --git a/modules/ocl/test/utility.hpp b/modules/ocl/test/utility.hpp index 1e17c6dbca..48c8bbcd9b 100644 --- a/modules/ocl/test/utility.hpp +++ b/modules/ocl/test/utility.hpp @@ -41,9 +41,15 @@ #ifndef __OPENCV_TEST_UTILITY_HPP__ #define __OPENCV_TEST_UTILITY_HPP__ + #define LOOP_TIMES 1 + #define MWIDTH 256 #define MHEIGHT 256 + +#define MIN_VALUE 171 +#define MAX_VALUE 351 + //#define RANDOMROI int randomInt(int minVal, int maxVal); double randomDouble(double minVal, double maxVal); @@ -73,6 +79,7 @@ double checkSimilarity(const cv::Mat &m1, const cv::Mat &m2); //oclMat create cv::ocl::oclMat createMat_ocl(cv::Size size, int type, bool useRoi = false); cv::ocl::oclMat loadMat_ocl(const cv::Mat& m, bool useRoi = false); + #define EXPECT_MAT_NORM(mat, eps) \ { \ EXPECT_LE(checkNorm(cv::Mat(mat)), eps) \ @@ -84,14 +91,7 @@ cv::ocl::oclMat loadMat_ocl(const cv::Mat& m, bool useRoi = false); ASSERT_EQ(mat1.size(), mat2.size()); \ EXPECT_LE(checkNorm(cv::Mat(mat1), cv::Mat(mat2)), eps); \ } -/* -#define EXPECT_MAT_NEAR(mat1, mat2, eps,s) \ -{ \ - ASSERT_EQ(mat1.type(), mat2.type()); \ - ASSERT_EQ(mat1.size(), mat2.size()); \ - EXPECT_LE(checkNorm(cv::Mat(mat1), cv::Mat(mat2)), eps)<