Merge pull request #1582 from sperrholz:ocl-arithm-additions

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

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

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 <cmpop> 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);

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);

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,

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 ////////////////////////////////////////////////////
 //////////////////////////////////////////////////////////////////////////////////////////////////////

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
+
     }
 }

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()));