diff --git a/modules/ocl/doc/matrix_reductions.rst b/modules/ocl/doc/matrix_reductions.rst index 4bedb944f..41161d8aa 100644 --- a/modules/ocl/doc/matrix_reductions.rst +++ b/modules/ocl/doc/matrix_reductions.rst @@ -23,6 +23,32 @@ Returns the number of non-zero elements in src Counts non-zero array elements. Supports all data types. +ocl::min +------------------ + +.. ocv:function:: void ocl::min(const oclMat &src1, const oclMat &src2, oclMat &dst) + + :param src1: the first input array. + + :param src2: the second input array, must be the same size and same type as ``src1``. + + :param dst: the destination array, it will have the same size and same type as ``src1``. + +Computes element-wise minima of two arrays. Supports all data types. + +ocl::max +------------------ + +.. ocv:function:: void ocl::max(const oclMat &src1, const oclMat &src2, oclMat &dst) + + :param src1: the first input array. + + :param src2: the second input array, must be the same size and same type as ``src1``. + + :param dst: the destination array, it will have the same size and same type as ``src1``. + +Computes element-wise maxima of two arrays. Supports all data types. + ocl::minMax ------------------ Returns void diff --git a/modules/ocl/doc/operations_on_matrices.rst b/modules/ocl/doc/operations_on_matrices.rst index 24a4ea1dc..7efd71967 100644 --- a/modules/ocl/doc/operations_on_matrices.rst +++ b/modules/ocl/doc/operations_on_matrices.rst @@ -3,6 +3,18 @@ Operations on Matrics .. highlight:: cpp +ocl::abs +------------------ +Returns void + +.. ocv:function:: void ocl::abs(const oclMat& src, oclMat& dst) + + :param src: input array. + + :param dst: destination array, it will have the same size and same type as ``src``. + +Computes per-element absolute values of the input array. Supports all data types. + ocl::absdiff ------------------ Returns void diff --git a/modules/ocl/include/opencv2/ocl/ocl.hpp b/modules/ocl/include/opencv2/ocl/ocl.hpp index 84e037238..58719f116 100644 --- a/modules/ocl/include/opencv2/ocl/ocl.hpp +++ b/modules/ocl/include/opencv2/ocl/ocl.hpp @@ -457,6 +457,14 @@ namespace cv // supports all data types CV_EXPORTS void divide(double scale, const oclMat &src1, oclMat &dst); + //! computes element-wise minimum of the two arrays (dst = min(src1, src2)) + // supports all data types + CV_EXPORTS void min(const oclMat &src1, const oclMat &src2, oclMat &dst); + + //! computes element-wise maximum of the two arrays (dst = max(src1, src2)) + // supports all data types + CV_EXPORTS void max(const oclMat &src1, const oclMat &src2, oclMat &dst); + //! compares elements of two arrays (dst = src1 src2) // supports all data types CV_EXPORTS void compare(const oclMat &src1, const oclMat &src2, oclMat &dst, int cmpop); @@ -465,6 +473,10 @@ namespace cv // supports all data types CV_EXPORTS void transpose(const oclMat &src, oclMat &dst); + //! computes element-wise absolute values of an array (dst = abs(src)) + // supports all data types + CV_EXPORTS void abs(const oclMat &src, oclMat &dst); + //! computes element-wise absolute difference of two arrays (dst = abs(src1 - src2)) // supports all data types CV_EXPORTS void absdiff(const oclMat &src1, const oclMat &src2, oclMat &dst); diff --git a/modules/ocl/src/arithm.cpp b/modules/ocl/src/arithm.cpp index 341f65c0a..301591269 100644 --- a/modules/ocl/src/arithm.cpp +++ b/modules/ocl/src/arithm.cpp @@ -57,10 +57,10 @@ using namespace cv; using namespace cv::ocl; ////////////////////////////////////////////////////////////////////////////// -/////////////////////// add subtract multiply divide ///////////////////////// +/////////////// add subtract multiply divide min max ///////////////////////// ////////////////////////////////////////////////////////////////////////////// -enum { ADD = 0, SUB, MUL, DIV, ABS_DIFF }; +enum { ADD = 0, SUB, MUL, DIV, ABS, ABS_DIFF, MIN, MAX }; static void arithmetic_run_generic(const oclMat &src1, const oclMat &src2, const Scalar & scalar, const oclMat & mask, oclMat &dst, int op_type, bool use_scalar = false) @@ -75,7 +75,7 @@ static void arithmetic_run_generic(const oclMat &src1, const oclMat &src2, const CV_Assert(src2.empty() || (!src2.empty() && src1.type() == src2.type() && src1.size() == src2.size())); CV_Assert(mask.empty() || (!mask.empty() && mask.type() == CV_8UC1 && mask.size() == src1.size())); - CV_Assert(op_type >= ADD && op_type <= ABS_DIFF); + CV_Assert(op_type >= ADD && op_type <= MAX); dst.create(src1.size(), src1.type()); @@ -93,7 +93,7 @@ static void arithmetic_run_generic(const oclMat &src1, const oclMat &src2, const const char * const typeMap[] = { "uchar", "char", "ushort", "short", "int", "float", "double" }; const char * const WTypeMap[] = { "short", "short", "int", "int", "int", "float", "double" }; - const char * const funcMap[] = { "FUNC_ADD", "FUNC_SUB", "FUNC_MUL", "FUNC_DIV", "FUNC_ABS_DIFF" }; + const char * const funcMap[] = { "FUNC_ADD", "FUNC_SUB", "FUNC_MUL", "FUNC_DIV", "FUNC_ABS", "FUNC_ABS_DIFF", "FUNC_MIN", "FUNC_MAX" }; const char * const channelMap[] = { "", "", "2", "4", "4" }; bool haveScalar = use_scalar || src2.empty(); @@ -205,10 +205,26 @@ void cv::ocl::divide(double scalar, const oclMat &src, oclMat &dst) arithmetic_run_generic(src, oclMat(), Scalar::all(scalar), oclMat(), dst, DIV); } +void cv::ocl::min(const oclMat &src1, const oclMat &src2, oclMat &dst) +{ + arithmetic_run_generic(src1, src2, Scalar::all(0), oclMat(), dst, MIN); +} + +void cv::ocl::max(const oclMat &src1, const oclMat &src2, oclMat &dst) +{ + arithmetic_run_generic(src1, src2, Scalar::all(0), oclMat(), dst, MAX); +} + ////////////////////////////////////////////////////////////////////////////// -///////////////////////////////// Absdiff //////////////////////////////////// +/////////////////////////////Abs, Absdiff //////////////////////////////////// ////////////////////////////////////////////////////////////////////////////// +void cv::ocl::abs(const oclMat &src, oclMat &dst) +{ + // explicitly uses use_scalar (even if zero) so that the correct kernel is used + arithmetic_run_generic(src, oclMat(), Scalar(), oclMat(), dst, ABS, true); +} + void cv::ocl::absdiff(const oclMat &src1, const oclMat &src2, oclMat &dst) { arithmetic_run_generic(src1, src2, Scalar(), oclMat(), dst, ABS_DIFF); diff --git a/modules/ocl/src/haar.cpp b/modules/ocl/src/haar.cpp index 05a76aa84..71777f148 100644 --- a/modules/ocl/src/haar.cpp +++ b/modules/ocl/src/haar.cpp @@ -926,7 +926,7 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS loopcount = 1; n_factors = 1; sizev.push_back(minSize); - scalev.push_back( min(cvRound(minSize.width / winsize0.width), cvRound(minSize.height / winsize0.height)) ); + scalev.push_back( std::min(cvRound(minSize.width / winsize0.width), cvRound(minSize.height / winsize0.height)) ); } detect_piramid_info *scaleinfo = (detect_piramid_info *)malloc(sizeof(detect_piramid_info) * loopcount); @@ -1555,7 +1555,7 @@ void cv::ocl::OclCascadeClassifierBuf::CreateFactorRelatedBufs( { loopcount = 1; sizev.push_back(minSize); - scalev.push_back( min(cvRound(minSize.width / winSize0.width), cvRound(minSize.height / winSize0.height)) ); + scalev.push_back( std::min(cvRound(minSize.width / winSize0.width), cvRound(minSize.height / winSize0.height)) ); } ((OclBuffers *)buffers)->pbuffer = openCLCreateBuffer(cv::ocl::Context::getContext(), CL_MEM_READ_ONLY, diff --git a/modules/ocl/src/opencl/arithm_add.cl b/modules/ocl/src/opencl/arithm_add.cl index 40caba5a9..cd9fae605 100644 --- a/modules/ocl/src/opencl/arithm_add.cl +++ b/modules/ocl/src/opencl/arithm_add.cl @@ -86,6 +86,14 @@ dst[dst_index] = convertToT(value); #endif +#if defined (FUNC_MIN) +#define EXPRESSION dst[dst_index] = min( src1[src1_index], src2[src2_index] ); +#endif + +#if defined (FUNC_MAX) +#define EXPRESSION dst[dst_index] = max( src1[src1_index], src2[src2_index] ); +#endif + ////////////////////////////////////////////////////////////////////////////////////////////////////// ///////////////////////////////////////////// ADD //////////////////////////////////////////////////// ////////////////////////////////////////////////////////////////////////////////////////////////////// diff --git a/modules/ocl/src/opencl/arithm_add_scalar.cl b/modules/ocl/src/opencl/arithm_add_scalar.cl index 4a0167fd5..f908b89f5 100644 --- a/modules/ocl/src/opencl/arithm_add_scalar.cl +++ b/modules/ocl/src/opencl/arithm_add_scalar.cl @@ -68,6 +68,12 @@ dst[dst_index] = src1[src1_index] == zero ? zero : convertToT(scalar[0] / convertToWT(src1[src1_index])); #endif +#if defined (FUNC_ABS) +#define EXPRESSION \ + T value = (src1[src1_index] > 0) ? src1[src1_index] : -src1[src1_index]; \ + dst[dst_index] = value; +#endif + #if defined (FUNC_ABS_DIFF) #define EXPRESSION WT value = convertToWT(src1[src1_index]) - scalar[0]; \ value = value > (WT)(0) ? value : -value; \ @@ -92,5 +98,6 @@ __kernel void arithm_binary_op_scalar (__global T *src1, int src1_step, int src1 int dst_index = mad24(y, dst_step, x + dst_offset); EXPRESSION + } } diff --git a/modules/ocl/test/test_arithm.cpp b/modules/ocl/test/test_arithm.cpp index 59565a887..db05f4556 100644 --- a/modules/ocl/test/test_arithm.cpp +++ b/modules/ocl/test/test_arithm.cpp @@ -453,6 +453,52 @@ TEST_P(Div, Mat_Scalar) //////////////////////////////// Absdiff ///////////////////////////////////////////////// +typedef ArithmTestBase Min; + +TEST_P(Min, Mat) +{ + for (int j = 0; j < LOOP_TIMES; j++) + { + random_roi(); + + dst1_roi = cv::min(src1_roi, src2_roi); + cv::ocl::min(gsrc1, gsrc2, gdst1); + Near(0); + } +} + +typedef ArithmTestBase Max; + +TEST_P(Max, Mat) +{ + for (int j = 0; j < LOOP_TIMES; j++) + { + random_roi(); + + dst1_roi = cv::min(src1_roi, src2_roi); + cv::ocl::min(gsrc1, gsrc2, gdst1); + Near(0); + } +} + +//////////////////////////////// Abs ///////////////////////////////////////////////////// + +typedef ArithmTestBase Abs; + +TEST_P(Abs, Abs) +{ + for (int j = 0; j < LOOP_TIMES; j++) + { + random_roi(); + + dst1_roi = cv::abs(src1_roi); + cv::ocl::abs(gsrc1, gdst1); + Near(0); + } +} + +//////////////////////////////// Absdiff ///////////////////////////////////////////////// + typedef ArithmTestBase Absdiff; TEST_P(Absdiff, Mat) @@ -1451,6 +1497,9 @@ INSTANTIATE_TEST_CASE_P(Arithm, Add, Combine(testing::Range(CV_8U, CV_USRTYPE1), INSTANTIATE_TEST_CASE_P(Arithm, Sub, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool())); INSTANTIATE_TEST_CASE_P(Arithm, Mul, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool())); INSTANTIATE_TEST_CASE_P(Arithm, Div, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool())); +INSTANTIATE_TEST_CASE_P(Arithm, Min, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool())); +INSTANTIATE_TEST_CASE_P(Arithm, Max, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool())); +INSTANTIATE_TEST_CASE_P(Arithm, Abs, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool())); INSTANTIATE_TEST_CASE_P(Arithm, Absdiff, Combine(testing::Range(CV_8U, CV_USRTYPE1), testing::Range(1, 5), Bool())); INSTANTIATE_TEST_CASE_P(Arithm, CartToPolar, Combine(Values(CV_32F, CV_64F), testing::Range(1, 5), Bool())); INSTANTIATE_TEST_CASE_P(Arithm, PolarToCart, Combine(Values(CV_32F, CV_64F), testing::Range(1, 5), Bool()));