Merge pull request #1756 from alalek:ocl_workaround_memory_leaks_with_subbuffer

This reverts commit d63a38e9bf.

Conflicts:
	modules/ocl/test/test_ml.cpp

.gitignore

@@ -1,10 +1,9 @@
-*.pyc
-.DS_Store
-refman.rst
-OpenCV4Tegra/
-tegra/
-*.user
-.sw[a-z]
-.*.swp
-tags
 *.autosave
+*.pyc
+*.user
+.*.swp
+.DS_Store
+.sw[a-z]
+/modules/refman.rst
+tags
+tegra/

doc/tutorials/introduction/android_binary_package/O4A_SDK.rst

@@ -51,7 +51,7 @@ The structure of package contents looks as follows:
     OpenCV-2.4.7-android-sdk
     |_ apk
     |   |_ OpenCV_2.4.7_binary_pack_armv7a.apk
-    |   |_ OpenCV_2.4.7_Manager_2.13_XXX.apk
+    |   |_ OpenCV_2.4.7_Manager_2.14_XXX.apk
     |
     |_ doc
     |_ samples
@@ -295,7 +295,7 @@ Well, running samples from Eclipse is very simple:
   .. code-block:: sh
     :linenos:
 
-    <Android SDK path>/platform-tools/adb install <OpenCV4Android SDK path>/apk/OpenCV_2.4.7_Manager_2.13_armv7a-neon.apk
+    <Android SDK path>/platform-tools/adb install <OpenCV4Android SDK path>/apk/OpenCV_2.4.7_Manager_2.14_armv7a-neon.apk
 
   .. note:: ``armeabi``, ``armv7a-neon``, ``arm7a-neon-android8``, ``mips`` and ``x86`` stand for
             platform targets:

modules/core/src/copy.cpp

@@ -485,6 +485,7 @@ void repeat(InputArray _src, int ny, int nx, OutputArray _dst)
 {
     Mat src = _src.getMat();
     CV_Assert( src.dims <= 2 );
+    CV_Assert( ny > 0 && nx > 0 );
 
     _dst.create(src.rows*ny, src.cols*nx, src.type());
     Mat dst = _dst.getMat();

modules/highgui/src/cap.cpp

@@ -523,7 +523,7 @@ bool VideoCapture::retrieve(Mat& image, int channel)
         return false;
     }
     if(_img->origin == IPL_ORIGIN_TL)
-        image = Mat(_img);
+        Mat(_img).copyTo(image);
     else
     {
         Mat temp(_img);

modules/imgproc/doc/filtering.rst

@@ -416,24 +416,23 @@ adaptiveBilateralFilter
 -----------------------
 Applies the adaptive bilateral filter to an image.
 
-.. ocv:function:: void adaptiveBilateralFilter( InputArray src, OutputArray dst, Size ksize, double sigmaSpace, Point anchor=Point(-1, -1), int borderType=BORDER_DEFAULT )
+.. ocv:function:: void adaptiveBilateralFilter( InputArray src, OutputArray dst, Size ksize, double sigmaSpace, double maxSigmaColor = 20.0, Point anchor=Point(-1, -1), int borderType=BORDER_DEFAULT )
 
 .. ocv:pyfunction:: cv2.adaptiveBilateralFilter(src, ksize, sigmaSpace[, dst[, anchor[, borderType]]]) -> dst
 
-    :param src: Source 8-bit, 1-channel or 3-channel image.
+    :param src: The source image
 
-    :param dst: Destination image of the same size and type as  ``src`` .
+    :param dst: The destination image; will have the same size and the same type as src
 
-    :param ksize: filter kernel size.
+    :param ksize: The kernel size. This is the neighborhood where the local variance will be calculated, and where pixels will contribute (in a weighted manner).
 
-    :param sigmaSpace: Filter sigma in the coordinate space. It has similar meaning with ``sigmaSpace`` in ``bilateralFilter``.
+    :param sigmaSpace: Filter sigma in the coordinate space. Larger value of the parameter means that farther pixels will influence each other (as long as their colors are close enough; see sigmaColor). Then d>0, it specifies the neighborhood size regardless of sigmaSpace, otherwise d is proportional to sigmaSpace.
 
-    :param anchor: anchor point; default value ``Point(-1,-1)`` means that the anchor is at the kernel center. Only default value is supported now.
+    :param maxSigmaColor: Maximum allowed sigma color (will clamp the value calculated in the ksize neighborhood. Larger value of the parameter means that more dissimilar pixels will influence each other (as long as their colors are close enough; see sigmaColor). Then d>0, it specifies the neighborhood size regardless of sigmaSpace, otherwise d is proportional to sigmaSpace.
 
-    :param borderType: border mode used to extrapolate pixels outside of the image.
-
-The function applies adaptive bilateral filtering to the input image. This filter is similar to ``bilateralFilter``, in that dissimilarity from and distance to the center pixel is punished. Instead of using ``sigmaColor``, we employ the variance of pixel values in the neighbourhood.
+    :param borderType: Pixel extrapolation method.
 
+A main part of our strategy will be to load each raw pixel once, and reuse it to calculate all pixels in the output (filtered) image that need this pixel value. The math of the filter is that of the usual bilateral filter, except that the sigma color is calculated in the neighborhood, and clamped by the optional input value.
 
 
 blur

modules/imgproc/include/opencv2/imgproc/imgproc.hpp

@@ -400,7 +400,7 @@ CV_EXPORTS_W void bilateralFilter( InputArray src, OutputArray dst, int d,
                                    int borderType=BORDER_DEFAULT );
 //! smooths the image using adaptive bilateral filter
 CV_EXPORTS_W void adaptiveBilateralFilter( InputArray src, OutputArray dst, Size ksize,
-                                           double sigmaSpace, Point anchor=Point(-1, -1),
+                                           double sigmaSpace, double maxSigmaColor = 20.0, Point anchor=Point(-1, -1),
                                            int borderType=BORDER_DEFAULT );
 //! smooths the image using the box filter. Each pixel is processed in O(1) time
 CV_EXPORTS_W void boxFilter( InputArray src, OutputArray dst, int ddepth,

modules/imgproc/src/smooth.cpp

@@ -2279,15 +2279,24 @@ void cv::bilateralFilter( InputArray _src, OutputArray _dst, int d,
 
 namespace cv
 {
-#define CALCVAR 1
-#define FIXED_WEIGHT 0
+#ifndef ABF_CALCVAR
+#define ABF_CALCVAR 1
+#endif
+
+#ifndef ABF_FIXED_WEIGHT
+#define ABF_FIXED_WEIGHT 0
+#endif
+
+#ifndef ABF_GAUSSIAN
+#define ABF_GAUSSIAN 1
+#endif
 
 class adaptiveBilateralFilter_8u_Invoker :
     public ParallelLoopBody
 {
 public:
-    adaptiveBilateralFilter_8u_Invoker(Mat& _dest, const Mat& _temp, Size _ksize, double _sigma_space, Point _anchor) :
-        temp(&_temp), dest(&_dest), ksize(_ksize), sigma_space(_sigma_space), anchor(_anchor)
+    adaptiveBilateralFilter_8u_Invoker(Mat& _dest, const Mat& _temp, Size _ksize, double _sigma_space, double _maxSigmaColor, Point _anchor) :
+        temp(&_temp), dest(&_dest), ksize(_ksize), sigma_space(_sigma_space), maxSigma_Color(_maxSigmaColor), anchor(_anchor)
     {
         if( sigma_space <= 0 )
             sigma_space = 1;
@@ -2300,7 +2309,11 @@ public:
         for(int y=-h; y<=h; y++)
             for(int x=-w; x<=w; x++)
         {
+#if ABF_GAUSSIAN
+            space_weight[idx++] = (float)exp ( -0.5*(x * x + y * y)/sigma2);
+#else
             space_weight[idx++] = (float)(sigma2 / (sigma2 + x * x + y * y));
+#endif
         }
     }
     virtual void operator()(const Range& range) const
@@ -2336,7 +2349,7 @@ public:
                     int startLMJ = 0;
                     int endLMJ  = ksize.width  - 1;
                     int howManyAll = (anX *2 +1)*(ksize.width );
-#if CALCVAR
+#if ABF_CALCVAR
                     for(int x = startLMJ; x< endLMJ; x++)
                     {
                         tptr = temp->ptr(startY + x) +j;
@@ -2348,8 +2361,14 @@ public:
                         }
                     }
                     var = ( (sumValSqr * howManyAll)- sumVal * sumVal )  /  ( (float)(howManyAll*howManyAll));
+
+                    if(var < 0.01)
+                        var = 0.01f;
+                    else if(var > (float)(maxSigma_Color*maxSigma_Color) )
+                        var =  (float)(maxSigma_Color*maxSigma_Color) ;
+
 #else
-                    var = 900.0;
+                    var = maxSigmaColor*maxSigmaColor;
 #endif
                     startLMJ = 0;
                     endLMJ = ksize.width;
@@ -2360,13 +2379,18 @@ public:
                         tptr = temp->ptr(startY + x) +j;
                         for(int y=-anX; y<=anX; y++)
                         {
-#if FIXED_WEIGHT
+#if ABF_FIXED_WEIGHT
                             weight = 1.0;
 #else
                             currVal = tptr[cn*(y+anX)];
                             currWRTCenter = currVal - currValCenter;
 
-                            weight = var / ( var + (currWRTCenter * currWRTCenter) ) * space_weight[x*ksize.width+y+anX];;
+#if ABF_GAUSSIAN
+                            weight = exp ( -0.5f * currWRTCenter * currWRTCenter/var ) * space_weight[x*ksize.width+y+anX];
+#else
+                            weight = var / ( var + (currWRTCenter * currWRTCenter) ) * space_weight[x*ksize.width+y+anX];
+#endif
+
 #endif
                             tmpSum += ((float)tptr[cn*(y+anX)] * weight);
                             totalWeight += weight;
@@ -2401,7 +2425,8 @@ public:
                     int startLMJ = 0;
                     int endLMJ  = ksize.width - 1;
                     int howManyAll = (anX *2 +1)*(ksize.width);
-#if CALCVAR
+#if ABF_CALCVAR
+                    float max_var = (float)( maxSigma_Color*maxSigma_Color);
                     for(int x = startLMJ; x< endLMJ; x++)
                     {
                         tptr = temp->ptr(startY + x) +j;
@@ -2416,11 +2441,27 @@ public:
                             sumValSqr_r += (currVal_r *currVal_r);
                         }
                     }
-                    var_b = ( (sumValSqr_b * howManyAll)- sumVal_b * sumVal_b )  /  ( (float)(howManyAll*howManyAll));
-                    var_g = ( (sumValSqr_g * howManyAll)- sumVal_g * sumVal_g )  /  ( (float)(howManyAll*howManyAll));
-                    var_r = ( (sumValSqr_r * howManyAll)- sumVal_r * sumVal_r )  /  ( (float)(howManyAll*howManyAll));
+                    var_b =  ( (sumValSqr_b * howManyAll)- sumVal_b * sumVal_b )  /  ( (float)(howManyAll*howManyAll));
+                    var_g =  ( (sumValSqr_g * howManyAll)- sumVal_g * sumVal_g )  /  ( (float)(howManyAll*howManyAll));
+                    var_r =  ( (sumValSqr_r * howManyAll)- sumVal_r * sumVal_r )  /  ( (float)(howManyAll*howManyAll));
+
+                    if(var_b < 0.01)
+                        var_b = 0.01f;
+                    else if(var_b > max_var )
+                        var_b =  (float)(max_var) ;
+
+                    if(var_g < 0.01)
+                        var_g = 0.01f;
+                    else if(var_g > max_var )
+                        var_g =  (float)(max_var) ;
+
+                    if(var_r < 0.01)
+                        var_r = 0.01f;
+                    else if(var_r > max_var )
+                        var_r =  (float)(max_var) ;
+
 #else
-                    var_b = 900.0; var_g = 900.0;var_r = 900.0;
+                    var_b = maxSigma_Color*maxSigma_Color; var_g = maxSigma_Color*maxSigma_Color; var_r = maxSigma_Color*maxSigma_Color;
 #endif
                     startLMJ = 0;
                     endLMJ = ksize.width;
@@ -2431,7 +2472,7 @@ public:
                         tptr = temp->ptr(startY + x) +j;
                         for(int y=-anX; y<=anX; y++)
                         {
-#if FIXED_WEIGHT
+#if ABF_FIXED_WEIGHT
                             weight_b = 1.0;
                             weight_g = 1.0;
                             weight_r = 1.0;
@@ -2442,9 +2483,16 @@ public:
                             currWRTCenter_r = currVal_r - currValCenter_r;
 
                             float cur_spw = space_weight[x*ksize.width+y+anX];
+
+#if ABF_GAUSSIAN
+                            weight_b = exp( -0.5f * currWRTCenter_b * currWRTCenter_b/ var_b ) * cur_spw;
+                            weight_g = exp( -0.5f * currWRTCenter_g * currWRTCenter_g/ var_g ) * cur_spw;
+                            weight_r = exp( -0.5f * currWRTCenter_r * currWRTCenter_r/ var_r ) * cur_spw;
+#else
                             weight_b = var_b / ( var_b + (currWRTCenter_b * currWRTCenter_b) ) * cur_spw;
                             weight_g = var_g / ( var_g + (currWRTCenter_g * currWRTCenter_g) ) * cur_spw;
                             weight_r = var_r / ( var_r + (currWRTCenter_r * currWRTCenter_r) ) * cur_spw;
+#endif
 #endif
                             tmpSum_b += ((float)tptr[cn*(y+anX)]   * weight_b);
                             tmpSum_g += ((float)tptr[cn*(y+anX)+1] * weight_g);
@@ -2468,10 +2516,11 @@ private:
     Mat *dest;
     Size ksize;
     double sigma_space;
+    double maxSigma_Color;
     Point anchor;
     vector<float> space_weight;
 };
-static void adaptiveBilateralFilter_8u( const Mat& src, Mat& dst, Size ksize, double sigmaSpace, Point anchor, int borderType )
+static void adaptiveBilateralFilter_8u( const Mat& src, Mat& dst, Size ksize, double sigmaSpace, double maxSigmaColor, Point anchor, int borderType )
 {
     Size size = src.size();
 
@@ -2481,12 +2530,12 @@ static void adaptiveBilateralFilter_8u( const Mat& src, Mat& dst, Size ksize, do
     Mat temp;
     copyMakeBorder(src, temp, anchor.x, anchor.y, anchor.x, anchor.y, borderType);
 
-    adaptiveBilateralFilter_8u_Invoker body(dst, temp, ksize, sigmaSpace, anchor);
+    adaptiveBilateralFilter_8u_Invoker body(dst, temp, ksize, sigmaSpace, maxSigmaColor, anchor);
     parallel_for_(Range(0, size.height), body, dst.total()/(double)(1<<16));
 }
 }
 void cv::adaptiveBilateralFilter( InputArray _src, OutputArray _dst, Size ksize,
-                                  double sigmaSpace, Point anchor, int borderType )
+                                  double sigmaSpace, double maxSigmaColor, Point anchor, int borderType )
 {
     Mat src = _src.getMat();
     _dst.create(src.size(), src.type());
@@ -2496,7 +2545,7 @@ void cv::adaptiveBilateralFilter( InputArray _src, OutputArray _dst, Size ksize,
 
     anchor = normalizeAnchor(anchor,ksize);
     if( src.depth() == CV_8U )
-        adaptiveBilateralFilter_8u( src, dst, ksize, sigmaSpace, anchor, borderType );
+        adaptiveBilateralFilter_8u( src, dst, ksize, sigmaSpace, maxSigmaColor, anchor, borderType );
     else
         CV_Error( CV_StsUnsupportedFormat,
         "Adaptive Bilateral filtering is only implemented for 8u images" );

modules/ocl/doc/camera_calibration_and_3D_reconstruction.rst

@@ -86,8 +86,6 @@ Enables the stereo correspondence operator that finds the disparity for the spec
 
     :param disparity: Output disparity map. It is a  ``CV_8UC1``  image with the same size as the input images.
 
-    :param stream: Stream for the asynchronous version.
-
 
 ocl::StereoBM_OCL::checkIfGpuCallReasonable
 -----------------------------------------------
@@ -218,8 +216,6 @@ Enables the stereo correspondence operator that finds the disparity for the spec
 
     :param disparity: Output disparity map. If  ``disparity``  is empty, the output type is  ``CV_16SC1`` . Otherwise, the type is retained.
 
-    :param stream: Stream for the asynchronous version.
-
 ocl::StereoConstantSpaceBP
 ------------------------------
 .. ocv:class:: ocl::StereoConstantSpaceBP
@@ -330,5 +326,3 @@ Enables the stereo correspondence operator that finds the disparity for the spec
     :param right: Right image with the same size and the same type as the left one.
 
     :param disparity: Output disparity map. If  ``disparity``  is empty, the output type is  ``CV_16SC1`` . Otherwise, the output type is  ``disparity.type()`` .
-
-    :param stream: Stream for the asynchronous version.

modules/ocl/doc/data_structures.rst

@@ -5,185 +5,193 @@ Data Structures
 
 OpenCV C++ 1-D or 2-D dense array class ::
 
-        class CV_EXPORTS oclMat
-        {
-        public:
-            //! default constructor
-            oclMat();
-            //! constructs oclMatrix of the specified size and type (_type is CV_8UC1, CV_64FC3, CV_32SC(12) etc.)
-            oclMat(int rows, int cols, int type);
-            oclMat(Size size, int type);
-            //! constucts oclMatrix and fills it with the specified value _s.
-            oclMat(int rows, int cols, int type, const Scalar &s);
-            oclMat(Size size, int type, const Scalar &s);
-            //! copy constructor
-            oclMat(const oclMat &m);
+    class CV_EXPORTS oclMat
+    {
+    public:
+        //! default constructor
+        oclMat();
+        //! constructs oclMatrix of the specified size and type (_type is CV_8UC1, CV_64FC3, CV_32SC(12) etc.)
+        oclMat(int rows, int cols, int type);
+        oclMat(Size size, int type);
+        //! constucts oclMatrix and fills it with the specified value _s.
+        oclMat(int rows, int cols, int type, const Scalar &s);
+        oclMat(Size size, int type, const Scalar &s);
+        //! copy constructor
+        oclMat(const oclMat &m);
 
-            //! constructor for oclMatrix headers pointing to user-allocated data
-            oclMat(int rows, int cols, int type, void *data, size_t step = Mat::AUTO_STEP);
-            oclMat(Size size, int type, void *data, size_t step = Mat::AUTO_STEP);
+        //! constructor for oclMatrix headers pointing to user-allocated data
+        oclMat(int rows, int cols, int type, void *data, size_t step = Mat::AUTO_STEP);
+        oclMat(Size size, int type, void *data, size_t step = Mat::AUTO_STEP);
 
-            //! creates a matrix header for a part of the bigger matrix
-            oclMat(const oclMat &m, const Range &rowRange, const Range &colRange);
-            oclMat(const oclMat &m, const Rect &roi);
+        //! creates a matrix header for a part of the bigger matrix
+        oclMat(const oclMat &m, const Range &rowRange, const Range &colRange);
+        oclMat(const oclMat &m, const Rect &roi);
 
-            //! builds oclMat from Mat. Perfom blocking upload to device.
-            explicit oclMat (const Mat &m);
+        //! builds oclMat from Mat. Perfom blocking upload to device.
+        explicit oclMat (const Mat &m);
 
-            //! destructor - calls release()
-            ~oclMat();
+        //! destructor - calls release()
+        ~oclMat();
 
-            //! assignment operators
-            oclMat &operator = (const oclMat &m);
-            //! assignment operator. Perfom blocking upload to device.
-            oclMat &operator = (const Mat &m);
+        //! assignment operators
+        oclMat &operator = (const oclMat &m);
+        //! assignment operator. Perfom blocking upload to device.
+        oclMat &operator = (const Mat &m);
+        oclMat &operator = (const oclMatExpr& expr);
+
+        //! pefroms blocking upload data to oclMat.
+        void upload(const cv::Mat &m);
 
 
-            //! pefroms blocking upload data to oclMat.
-            void upload(const cv::Mat &m);
+        //! downloads data from device to host memory. Blocking calls.
+        operator Mat() const;
+        void download(cv::Mat &m) const;
 
+        //! convert to _InputArray
+        operator _InputArray();
 
-            //! downloads data from device to host memory. Blocking calls.
-            operator Mat() const;
-            void download(cv::Mat &m) const;
+        //! convert to _OutputArray
+        operator _OutputArray();
 
+        //! returns a new oclMatrix header for the specified row
+        oclMat row(int y) const;
+        //! returns a new oclMatrix header for the specified column
+        oclMat col(int x) const;
+        //! ... for the specified row span
+        oclMat rowRange(int startrow, int endrow) const;
+        oclMat rowRange(const Range &r) const;
+        //! ... for the specified column span
+        oclMat colRange(int startcol, int endcol) const;
+        oclMat colRange(const Range &r) const;
 
-            //! returns a new oclMatrix header for the specified row
-            oclMat row(int y) const;
-            //! returns a new oclMatrix header for the specified column
-            oclMat col(int x) const;
-            //! ... for the specified row span
-            oclMat rowRange(int startrow, int endrow) const;
-            oclMat rowRange(const Range &r) const;
-            //! ... for the specified column span
-            oclMat colRange(int startcol, int endcol) const;
-            oclMat colRange(const Range &r) const;
+        //! returns deep copy of the oclMatrix, i.e. the data is copied
+        oclMat clone() const;
 
-            //! returns deep copy of the oclMatrix, i.e. the data is copied
-            oclMat clone() const;
-            //! copies the oclMatrix content to "m".
-            // It calls m.create(this->size(), this->type()).
-            // It supports any data type
-            void copyTo( oclMat &m ) const;
-            //! copies those oclMatrix elements to "m" that are marked with non-zero mask elements.
-            //It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
-            void copyTo( oclMat &m, const oclMat &mask ) const;
-            //! converts oclMatrix to another datatype with optional scalng. See cvConvertScale.
-            //It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
-            void convertTo( oclMat &m, int rtype, double alpha = 1, double beta = 0 ) const;
+        //! copies those oclMatrix elements to "m" that are marked with non-zero mask elements.
+        // It calls m.create(this->size(), this->type()).
+        // It supports any data type
+        void copyTo( oclMat &m, const oclMat &mask = oclMat()) const;
 
-            void assignTo( oclMat &m, int type = -1 ) const;
+        //! converts oclMatrix to another datatype with optional scalng. See cvConvertScale.
+        void convertTo( oclMat &m, int rtype, double alpha = 1, double beta = 0 ) const;
 
-            //! sets every oclMatrix element to s
-            //It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
-            oclMat &operator = (const Scalar &s);
-            //! sets some of the oclMatrix elements to s, according to the mask
-            //It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
-            oclMat &setTo(const Scalar &s, const oclMat &mask = oclMat());
-            //! creates alternative oclMatrix header for the same data, with different
-            // number of channels and/or different number of rows. see cvReshape.
-            oclMat reshape(int cn, int rows = 0) const;
+        void assignTo( oclMat &m, int type = -1 ) const;
 
-            //! allocates new oclMatrix data unless the oclMatrix already has specified size and type.
-            // previous data is unreferenced if needed.
-            void create(int rows, int cols, int type);
-            void create(Size size, int type);
-            //! decreases reference counter;
-            // deallocate the data when reference counter reaches 0.
-            void release();
+        //! sets every oclMatrix element to s
+        oclMat& operator = (const Scalar &s);
+        //! sets some of the oclMatrix elements to s, according to the mask
+        oclMat& setTo(const Scalar &s, const oclMat &mask = oclMat());
+        //! creates alternative oclMatrix header for the same data, with different
+        // number of channels and/or different number of rows. see cvReshape.
+        oclMat reshape(int cn, int rows = 0) const;
 
-            //! swaps with other smart pointer
-            void swap(oclMat &mat);
+        //! allocates new oclMatrix data unless the oclMatrix already has specified size and type.
+        // previous data is unreferenced if needed.
+        void create(int rows, int cols, int type);
+        void create(Size size, int type);
 
-            //! locates oclMatrix header within a parent oclMatrix. See below
-            void locateROI( Size &wholeSize, Point &ofs ) const;
-            //! moves/resizes the current oclMatrix ROI inside the parent oclMatrix.
-            oclMat &adjustROI( int dtop, int dbottom, int dleft, int dright );
-            //! extracts a rectangular sub-oclMatrix
-            // (this is a generalized form of row, rowRange etc.)
-            oclMat operator()( Range rowRange, Range colRange ) const;
-            oclMat operator()( const Rect &roi ) const;
+        //! allocates new oclMatrix with specified device memory type.
+        void createEx(int rows, int cols, int type, DevMemRW rw_type, DevMemType mem_type);
+        void createEx(Size size, int type, DevMemRW rw_type, DevMemType mem_type);
 
-            //! returns true if the oclMatrix data is continuous
-            // (i.e. when there are no gaps between successive rows).
-            // similar to CV_IS_oclMat_CONT(cvoclMat->type)
-            bool isContinuous() const;
-            //! returns element size in bytes,
-            // similar to CV_ELEM_SIZE(cvMat->type)
-            size_t elemSize() const;
-            //! returns the size of element channel in bytes.
-            size_t elemSize1() const;
-            //! returns element type, similar to CV_MAT_TYPE(cvMat->type)
-            int type() const;
-            //! returns element type, i.e. 8UC3 returns 8UC4 because in ocl
-            //! 3 channels element actually use 4 channel space
-            int ocltype() const;
-            //! returns element type, similar to CV_MAT_DEPTH(cvMat->type)
-            int depth() const;
-            //! returns element type, similar to CV_MAT_CN(cvMat->type)
-            int channels() const;
-            //! returns element type, return 4 for 3 channels element,
-            //!becuase 3 channels element actually use 4 channel space
-            int oclchannels() const;
-            //! returns step/elemSize1()
-            size_t step1() const;
-            //! returns oclMatrix size:
-            // width == number of columns, height == number of rows
-            Size size() const;
-            //! returns true if oclMatrix data is NULL
-            bool empty() const;
+        //! decreases reference counter;
+        // deallocate the data when reference counter reaches 0.
+        void release();
 
-            //! returns pointer to y-th row
-            uchar *ptr(int y = 0);
-            const uchar *ptr(int y = 0) const;
+        //! swaps with other smart pointer
+        void swap(oclMat &mat);
 
-            //! template version of the above method
-            template<typename _Tp> _Tp *ptr(int y = 0);
-            template<typename _Tp> const _Tp *ptr(int y = 0) const;
+        //! locates oclMatrix header within a parent oclMatrix. See below
+        void locateROI( Size &wholeSize, Point &ofs ) const;
+        //! moves/resizes the current oclMatrix ROI inside the parent oclMatrix.
+        oclMat& adjustROI( int dtop, int dbottom, int dleft, int dright );
+        //! extracts a rectangular sub-oclMatrix
+        // (this is a generalized form of row, rowRange etc.)
+        oclMat operator()( Range rowRange, Range colRange ) const;
+        oclMat operator()( const Rect &roi ) const;
 
-            //! matrix transposition
-            oclMat t() const;
+        oclMat& operator+=( const oclMat& m );
+        oclMat& operator-=( const oclMat& m );
+        oclMat& operator*=( const oclMat& m );
+        oclMat& operator/=( const oclMat& m );
 
-            /*! includes several bit-fields:
-              - the magic signature
-              - continuity flag
-              - depth
-              - number of channels
-              */
-            int flags;
-            //! the number of rows and columns
-            int rows, cols;
-            //! a distance between successive rows in bytes; includes the gap if any
-            size_t step;
-            //! pointer to the data(OCL memory object)
-            uchar *data;
+        //! returns true if the oclMatrix data is continuous
+        // (i.e. when there are no gaps between successive rows).
+        // similar to CV_IS_oclMat_CONT(cvoclMat->type)
+        bool isContinuous() const;
+        //! returns element size in bytes,
+        // similar to CV_ELEM_SIZE(cvMat->type)
+        size_t elemSize() const;
+        //! returns the size of element channel in bytes.
+        size_t elemSize1() const;
+        //! returns element type, similar to CV_MAT_TYPE(cvMat->type)
+        int type() const;
+        //! returns element type, i.e. 8UC3 returns 8UC4 because in ocl
+        //! 3 channels element actually use 4 channel space
+        int ocltype() const;
+        //! returns element type, similar to CV_MAT_DEPTH(cvMat->type)
+        int depth() const;
+        //! returns element type, similar to CV_MAT_CN(cvMat->type)
+        int channels() const;
+        //! returns element type, return 4 for 3 channels element,
+        //!becuase 3 channels element actually use 4 channel space
+        int oclchannels() const;
+        //! returns step/elemSize1()
+        size_t step1() const;
+        //! returns oclMatrix size:
+        // width == number of columns, height == number of rows
+        Size size() const;
+        //! returns true if oclMatrix data is NULL
+        bool empty() const;
 
-            //! pointer to the reference counter;
-            // when oclMatrix points to user-allocated data, the pointer is NULL
-            int *refcount;
+        //! returns pointer to y-th row
+        uchar* ptr(int y = 0);
+        const uchar *ptr(int y = 0) const;
 
-            //! helper fields used in locateROI and adjustROI
-            //datastart and dataend are not used in current version
-            uchar *datastart;
-            uchar *dataend;
+        //! template version of the above method
+        template<typename _Tp> _Tp *ptr(int y = 0);
+        template<typename _Tp> const _Tp *ptr(int y = 0) const;
 
-            //! OpenCL context associated with the oclMat object.
-            Context *clCxt;
-            //add offset for handle ROI, calculated in byte
-            int offset;
-            //add wholerows and wholecols for the whole matrix, datastart and dataend are no longer used
-            int wholerows;
-            int wholecols;
-        };
+        //! matrix transposition
+        oclMat t() const;
 
-Basically speaking, the oclMat is the mirror of Mat with the extension of ocl feature, the members have the same meaning and useage of Mat except following:
+        /*! includes several bit-fields:
+          - the magic signature
+          - continuity flag
+          - depth
+          - number of channels
+          */
+        int flags;
+        //! the number of rows and columns
+        int rows, cols;
+        //! a distance between successive rows in bytes; includes the gap if any
+        size_t step;
+        //! pointer to the data(OCL memory object)
+        uchar *data;
 
-datastart and dataend are replaced with wholerows and wholecols
+        //! pointer to the reference counter;
+        // when oclMatrix points to user-allocated data, the pointer is NULL
+        int *refcount;
 
-add clCxt for oclMat
+        //! helper fields used in locateROI and adjustROI
+        //datastart and dataend are not used in current version
+        uchar *datastart;
+        uchar *dataend;
 
-Only basic flags are supported in oclMat(i.e. depth number of channels)
+        //! OpenCL context associated with the oclMat object.
+        Context *clCxt;
+        //add offset for handle ROI, calculated in byte
+        int offset;
+        //add wholerows and wholecols for the whole matrix, datastart and dataend are no longer used
+        int wholerows;
+        int wholecols;
+    };
 
-All the 3-channel matrix(i.e. RGB image) are represented by 4-channel matrix in oclMat. It means 3-channel image have 4-channel space with the last channel unused. We provide a transparent interface to handle the difference between OpenCV Mat and oclMat.
+Basically speaking, the ``oclMat`` is the mirror of ``Mat`` with the extension of OCL feature, the members have the same meaning and useage of ``Mat`` except following:
 
-For example: If a oclMat has 3 channels, channels() returns 3 and oclchannels() returns 4
+* ``datastart`` and ``dataend`` are replaced with ``wholerows`` and ``wholecols``
+
+* Only basic flags are supported in ``oclMat`` (i.e. depth number of channels)
+
+* All the 3-channel matrix (i.e. RGB image) are represented by 4-channel matrix in ``oclMat``. It means 3-channel image have 4-channel space with the last channel unused. We provide a transparent interface to handle the difference between OpenCV ``Mat`` and ``oclMat``.
+    For example: If a ``oclMat`` has 3 channels, ``channels()`` returns 3 and ``oclchannels()`` returns 4

modules/ocl/doc/image_filtering.rst

@@ -497,23 +497,21 @@ ocl::adaptiveBilateralFilter
 --------------------------------
 Returns void
 
-.. ocv:function:: void ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, Point anchor = Point(-1, -1), int borderType=BORDER_DEFAULT)
+.. ocv:function:: void ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, double maxSigmaColor = 20.0, Point anchor = Point(-1, -1), int borderType=BORDER_DEFAULT)
 
     :param src: The source image
 
     :param dst: The destination image; will have the same size and the same type as src
 
-    :param ksize: The kernel size
+    :param ksize: The kernel size. This is the neighborhood where the local variance will be calculated, and where pixels will contribute (in a weighted manner).
 
     :param sigmaSpace: Filter sigma in the coordinate space. Larger value of the parameter means that farther pixels will influence each other (as long as their colors are close enough; see sigmaColor). Then d>0, it specifies the neighborhood size regardless of sigmaSpace, otherwise d is proportional to sigmaSpace.
 
+    :param maxSigmaColor: Maximum allowed sigma color (will clamp the value calculated in the ksize neighborhood. Larger value of the parameter means that more dissimilar pixels will influence each other (as long as their colors are close enough; see sigmaColor). Then d>0, it specifies the neighborhood size regardless of sigmaSpace, otherwise d is proportional to sigmaSpace.
+
     :param borderType: Pixel extrapolation method.
 
-A main part of our strategy will be to load each raw pixel once, and reuse it to calculate all pixels in the output (filtered) image that need this pixel value.
-
-.. math::
-
-    \emph{O}_i = \frac{1}{W_i}\sum\limits_{j\in{N(i)}}{\frac{1}{1+\frac{(V_i-V_j)^2}{\sigma_{N{'}(i)}^2}}*\frac{1}{1+\frac{d(i,j)^2}{\sum^2}}}V_j
+A main part of our strategy will be to load each raw pixel once, and reuse it to calculate all pixels in the output (filtered) image that need this pixel value. The math of the filter is that of the usual bilateral filter, except that the sigma color is calculated in the neighborhood, and clamped by the optional input value.
 
 Local memory organization
 

modules/ocl/doc/image_processing.rst

@@ -146,7 +146,7 @@ Returns void
 
 .. ocv:function:: void ocl::remap(const oclMat &src, oclMat &dst, oclMat &map1, oclMat &map2, int interpolation, int bordertype, const Scalar &value = Scalar())
 
-    :param src: Source image. Only CV_8UC1 and CV_32FC1 images are supported now.
+    :param src: Source image.
 
     :param dst: Destination image containing cornerness values. It has the same size as src and CV_32FC1 type.
 
@@ -156,11 +156,11 @@ Returns void
 
     :param interpolation: The interpolation method
 
-    :param bordertype: Pixel extrapolation method. Only BORDER_CONSTANT are supported now.
+    :param bordertype: Pixel extrapolation method.
 
     :param value: The border value if borderType==BORDER CONSTANT
 
-The function remap transforms the source image using the specified map: dst (x ,y) = src (map1(x , y) , map2(x , y)) where values of pixels with non-integer coordinates are computed using one of available interpolation methods. map1 and map2 can be encoded as separate floating-point maps in map1 and map2 respectively, or interleaved floating-point maps of (x,y) in map1. Supports CV_8UC1, CV_8UC3, CV_8UC4, CV_32FC1 , CV_32FC3 and CV_32FC4 data types.
+The function remap transforms the source image using the specified map: dst (x ,y) = src (map1(x , y) , map2(x , y)) where values of pixels with non-integer coordinates are computed using one of available interpolation methods. map1 and map2 can be encoded as separate floating-point maps in map1 and map2 respectively, or interleaved floating-point maps of (x,y) in map1.
 
 ocl::resize
 ------------------
@@ -222,7 +222,7 @@ ocl::cvtColor
 ------------------
 Returns void
 
-.. ocv:function:: void ocl::cvtColor(const oclMat &src, oclMat &dst, int code , int dcn = 0)
+.. ocv:function:: void ocl::cvtColor(const oclMat &src, oclMat &dst, int code, int dcn = 0)
 
     :param src: Source image.
 
@@ -250,7 +250,7 @@ Returns Threshold value
 
     :param type: Thresholding type
 
-The function applies fixed-level thresholding to a single-channel array. The function is typically used to get a bi-level (binary) image out of a grayscale image or for removing a noise, i.e. filtering out pixels with too small or too large values. There are several types of thresholding that the function supports that are determined by thresholdType. Supports only CV_32FC1 and CV_8UC1 data type.
+The function applies fixed-level thresholding to a single-channel array. The function is typically used to get a bi-level (binary) image out of a grayscale image or for removing a noise, i.e. filtering out pixels with too small or too large values. There are several types of thresholding that the function supports that are determined by thresholdType.
 
 ocl::buildWarpPlaneMaps
 -----------------------
@@ -311,4 +311,4 @@ Builds transformation maps for affine transformation.
 
     :param ymap: Y values with  ``CV_32FC1`` type.
 
-.. seealso:: :ocv:func:`ocl::warpAffine` , :ocv:func:`ocl::remap`
+.. seealso:: :ocv:func:`ocl::warpAffine` , :ocv:func:`ocl::remap`

modules/ocl/doc/introduction.rst

@@ -6,53 +6,68 @@ OpenCL Module Introduction
 General Information
 -------------------
 
-The OpenCV OCL module contains  a set of classes and functions that implement and accelerate select openCV functionality on OpenCL compatible devices. OpenCL is a Khronos standard, implemented by a variety of devices (CPUs, GPUs, FPGAs, ARM), abstracting the exact hardware details, while enabling vendors to provide native implementation for maximal acceleration on their hardware. The standard enjoys wide industry support, and the end user of the module will enjoy the data parallelism benefits that the specific platform/hardware may be capable of, in a platform/hardware independent manner.
+The OpenCV OCL module contains a set of classes and functions that implement and accelerate OpenCV functionality on OpenCL compatible devices. OpenCL is a Khronos standard, implemented by a variety of devices (CPUs, GPUs, FPGAs, ARM), abstracting the exact hardware details, while enabling vendors to provide native implementation for maximal acceleration on their hardware. The standard enjoys wide industry support, and the end user of the module will enjoy the data parallelism benefits that the specific platform/hardware may be capable of, in a platform/hardware independent manner.
 
-While in the future we hope to validate (and enable) the OCL module in all OpenCL capable devices, we currently develop and test on GPU devices only. This includes both discrete GPUs (NVidia, AMD), as well as integrated chips(AMD APU and intel HD devices). Performance of any particular algorithm will depend on the particular platform characteristics and capabilities. However, currently (as of 2.4.4), accuracy and  mathematical correctness has been verified to be identical to that of the pure CPU implementation on all tested GPU devices and platforms (both windows and linux).
+While in the future we hope to validate (and enable) the OCL module in all OpenCL capable devices, we currently develop and test on GPU devices only. This includes both discrete GPUs (NVidia, AMD), as well as integrated chips (AMD APU and Intel HD devices). Performance of any particular algorithm will depend on the particular platform characteristics and capabilities. However, currently, accuracy and  mathematical correctness has been verified to be identical to that of the pure CPU implementation on all tested GPU devices and platforms (both Windows and Linux).
 
 
-The OpenCV OCL module includes utility functions, low-level vision primitives, and high-level algorithms. The utility functions and low-level primitives provide a powerful infrastructure for developing fast vision algorithms taking advangtage of OCL whereas the high-level functionality (samples)includes some state-of-the-art algorithms (including LK Optical flow, and Face detection) ready to be used by the application developers. The module is also accompanied by an extensive performance and accuracy test suite.
+The OpenCV OCL module includes utility functions, low-level vision primitives, and high-level algorithms. The utility functions and low-level primitives provide a powerful infrastructure for developing fast vision algorithms taking advantage of OCL, whereas the high-level functionality (samples) includes some state-of-the-art algorithms (including LK Optical flow, and Face detection) ready to be used by the application developers. The module is also accompanied by an extensive performance and accuracy test suite.
 
-The OpenCV OCL module is designed for ease of use and does not require any knowledge of OpenCL. At a minimuml level, it can be viewed as a set of accelerators, that can take advantage of the high compute throughput that GPU/APU devices can provide. However, it can also be viewed as a starting point to really integratethe  built-in functionality with your own custom OpenCL kernels, with or without modifying the source of OpenCV-OCL. Of course, knowledge of OpenCL will certainly help, however we hope that OpenCV-OCL module, and  the kernels it contains in source code, can be very useful as a means of actually learning openCL. Such a knowledge would be necessary to further fine-tune any of the existing OpenCL kernels, or for extending the framework with new kernels. As of OpenCV 2.4.4, we introduce interoperability with OpenCL, enabling easy use of custom OpenCL kernels within the OpenCV framework.
+The OpenCV OCL module is designed for ease of use and does not require any knowledge of OpenCL. At a minimum level, it can be viewed as a set of accelerators, that can take advantage of the high compute throughput that GPU/APU devices can provide. However, it can also be viewed as a starting point to really integrate the built-in functionality with your own custom OpenCL kernels, with or without modifying the source of OpenCV-OCL. Of course, knowledge of OpenCL will certainly help, however we hope that OpenCV-OCL module, and the kernels it contains in source code, can be very useful as a means of actually learning openCL. Such a knowledge would be necessary to further fine-tune any of the existing OpenCL kernels, or for extending the framework with new kernels. As of OpenCV 2.4.4, we introduce interoperability with OpenCL, enabling easy use of custom OpenCL kernels within the OpenCV framework.
 
-To use the OCL module, you need to make sure that you have the OpenCL SDK provided with your device vendor. To correctly run the OCL module, you need to have the OpenCL runtime provide by the device vendor, typically the device driver.
+To correctly run the OCL module, you need to have the OpenCL runtime provided by the device vendor, typically the device driver.
 
-To enable OCL support, configure OpenCV using CMake with WITH\_OPENCL=ON. When the flag is set and if OpenCL SDK is installed, the full-featured OpenCV OCL module is built. Otherwise, the module may be not built. If you have AMD'S FFT and BLAS library, you can select it with WITH\_OPENCLAMDFFT=ON, WITH\_OPENCLAMDBLAS=ON.
+To enable OCL support, configure OpenCV using CMake with ``WITH_OPENCL=ON``. When the flag is set and if OpenCL SDK is installed, the full-featured OpenCV OCL module is built. Otherwise, the module may be not built. If you have AMD'S FFT and BLAS library, you can select it with ``WITH_OPENCLAMDFFT=ON``, ``WITH_OPENCLAMDBLAS=ON``.
 
-The ocl module can be found under the "modules" directory. In "modules/ocl/src" you can find the source code for the cpp class that wrap around the direct kernel invocation. The kernels themselves can be found in "modules/ocl/src/kernels."  Samples can be found under "samples/ocl." Accuracy tests can be found in "modules/ocl/test," and performance tests under "module/ocl/perf."
+The ocl module can be found under the "modules" directory. In "modules/ocl/src" you can find the source code for the cpp class that wrap around the direct kernel invocation. The kernels themselves can be found in "modules/ocl/src/opencl".  Samples can be found under "samples/ocl". Accuracy tests can be found in "modules/ocl/test", and performance tests under "module/ocl/perf".
 
 
 
+Right now, the user can select OpenCL device by specifying the environment variable ``OPENCV_OPENCL_DEVICE``. Variable format:
 
-Right now, the user should define the cv::ocl::Info class in the application and call cv::ocl::getDevice before any cv::ocl::func. This operation initialize OpenCL runtime and set the first found device as computing device. If there are more than one device and you want to use undefault device, you can call cv::ocl::setDevice then.
+.. code-block:: cpp
 
-In the current version, all the thread share the same context and device so the multi-devices are not supported. We will add this feature soon. If a function support 4-channel operator, it should support 3-channel operator as well, because All the 3-channel matrix(i.e. RGB image) are represented by 4-channel matrix in oclMat. It means 3-channel image have 4-channel space with the last channel unused. We provide a transparent interface to handle the difference between OpenCV Mat and oclMat.
+    <Platform>:<CPU|GPU|ACCELERATOR|nothing=GPU/CPU>:<DeviceName or ID>
+
+**Note:** Device ID range is: 0..9 (only one digit, 10 - it is a part of name)
+
+Samples:
+
+.. code-block:: cpp
+
+    '' = ':' = '::' = ':GPU|CPU:'
+    'AMD:GPU|CPU:'
+    'AMD::Tahiti'
+    ':GPU:1'
+    ':CPU:2'
+
+Also the user can use ``cv::ocl::setDevice`` function (with ``cv::ocl::getOpenCLPlatforms`` and ``cv::ocl::getOpenCLDevices``). This function initializes OpenCL runtime and setup the passed device as computing device.
+
+In the current version, all the thread share the same context and device so the multi-devices are not supported. We will add this feature soon. If a function support 4-channel operator, it should support 3-channel operator as well, because All the 3-channel matrix(i.e. RGB image) are represented by 4-channel matrix in ``oclMat``. It means 3-channel image have 4-channel space with the last channel unused. We provide a transparent interface to handle the difference between OpenCV Mat and ``oclMat``.
 
 Developer Notes
 -------------------
 
-In a heterogeneous device environment, there may be cost associated with data transfer. This would be the case, for example, when data needs to be moved from host memory (accessible to the CPU), to device memory (accessible to a discrete GPU). in the case of integrated graphics chips, there may be performance issues, relating to memory coherency between access from the GPU "part" of the integrated device, or the CPU "part." For best performance, in either case, it is recommended that you do not introduce dat transfers between CPU and the discrete GPU, except in the beginning and the end of the algorithmic pipeline.
+In a heterogeneous device environment, there may be cost associated with data transfer. This would be the case, for example, when data needs to be moved from host memory (accessible to the CPU), to device memory (accessible to a discrete GPU). in the case of integrated graphics chips, there may be performance issues, relating to memory coherency between access from the GPU "part" of the integrated device, or the CPU "part." For best performance, in either case, it is recommended that you do not introduce data transfers between CPU and the discrete GPU, except in the beginning and the end of the algorithmic pipeline.
 
 Some tidbits:
 
 1. OpenCL version should be larger than 1.1 with FULL PROFILE.
 
-2. Currently (2.4.4) the user call the cv::ocl::getDevice before any other function in the ocl module. This will initialize the OpenCL runtime and set the first found device as computing device. If there are more than one device and you want to use undefault device, you can call cv::ocl::setDevice thereafter.
-
 2. Currently there's only one OpenCL context and command queue. We hope to implement multi device and multi queue support in the future.
 
 3. Many kernels use 256 as its workgroup size if possible, so the max work group size of the device must larger than 256. All GPU devices we are aware of indeed support 256 workitems in a workgroup, however non GPU devices may not. This will be improved in the future.
 
-4. If the device does not support double arithetic, we revert to float.
+4. If the device does not support double arithmetic, then functions' implementation generates an error.
 
-5. The oclMat uses buffer object, not image object.
+5. The ``oclMat`` uses buffer object, not image object.
 
-6. All the 3-channel matrices(i.e. RGB image) are represented by 4-channel matrices in oclMat, with the last channel unused. We provide a transparent interface to handle the difference between OpenCV Mat and oclMat.
+6. All the 3-channel matrices (i.e. RGB image) are represented by 4-channel matrices in ``oclMat``, with the last channel unused. We provide a transparent interface to handle the difference between OpenCV Mat and ``oclMat``.
 
-7. All the matrix in oclMat is aligned in column(now the alignment factor is 32 byte). It means, if a matrix is n columns m rows with the element size x byte, we will assign ALIGNMENT(x*n) bytes for each column with the last ALIGNMENT(x*n) - x*n bytes unused, so there's small holes after each column if its size is not the multiply of ALIGN.
+7. All the matrix in ``oclMat`` is aligned in column (now the alignment factor for ``step`` is 32+ byte). It means, m.cols * m.elemSize() <= m.step.
 
-8. Data transfer between Mat and oclMat: If the CPU matrix is aligned in column, we will use faster API to transfer between Mat and oclMat, otherwise, we will use clEnqueueRead/WriteBufferRect to transfer data to guarantee the alignment. 3-channel matrix is an exception, it's directly transferred to a temp buffer and then padded to 4-channel matrix(also aligned) when uploading and do the reverse operation when downloading.
+8. Data transfer between Mat and ``oclMat``: If the CPU matrix is aligned in column, we will use faster API to transfer between Mat and ``oclMat``, otherwise, we will use clEnqueueRead/WriteBufferRect to transfer data to guarantee the alignment. 3-channel matrix is an exception, it's directly transferred to a temp buffer and then padded to 4-channel matrix(also aligned) when uploading and do the reverse operation when downloading.
 
-9. Data transfer between Mat and oclMat: ROI is a feature of OpenCV, which allow users process a sub rectangle of a matrix. When a CPU matrix which has ROI will be transfered to GPU, the whole matrix will be transfered and set ROI as CPU's. In a word, we always transfer the whole matrix despite whether it has ROI or not.
+9. Data transfer between Mat and ``oclMat``: ROI is a feature of OpenCV, which allow users process a sub rectangle of a matrix. When a CPU matrix which has ROI will be transfered to GPU, the whole matrix will be transfered and set ROI as CPU's. In a word, we always transfer the whole matrix despite whether it has ROI or not.
 
-10. All the kernel file should locate in ocl/src/kernels/ with the extension ".cl". ALL the kernel files are transformed to pure characters at compilation time in kernels.cpp, and the file name without extension is the name of the characters.
+10. All the kernel file should locate in "modules/ocl/src/opencl/" with the extension ".cl". All the kernel files are transformed to pure characters at compilation time in opencl_kernels.cpp, and the file name without extension is the name of the program sources.

modules/ocl/doc/video_analysis.rst

@@ -117,7 +117,6 @@ Computes a dense optical flow using the Gunnar Farneback's algorithm.
     :param frame1: Second 8-bit gray-scale input image
     :param flowx: Flow horizontal component
     :param flowy: Flow vertical component
-    :param s: Stream
 
 .. seealso:: :ocv:func:`calcOpticalFlowFarneback`
 
@@ -230,8 +229,6 @@ Interpolates frames (images) using provided optical flow (displacement field).
 
     :param buf: Temporary buffer, will have width x 6*height size, CV_32FC1 type and contain 6 oclMat: occlusion masks for first frame, occlusion masks for second, interpolated forward horizontal flow, interpolated forward vertical flow, interpolated backward horizontal flow, interpolated backward vertical flow.
 
-    :param stream: Stream for the asynchronous version.
-
 ocl::KalmanFilter
 --------------------
 .. ocv:class:: ocl::KalmanFilter
@@ -418,8 +415,6 @@ Updates the background model and returns the foreground mask.
 
     :param fgmask: The output foreground mask as an 8-bit binary image.
 
-    :param stream: Stream for the asynchronous version.
-
 
 ocl::MOG::getBackgroundImage
 --------------------------------
@@ -429,8 +424,6 @@ Computes a background image.
 
     :param backgroundImage: The output background image.
 
-    :param stream: Stream for the asynchronous version.
-
 
 ocl::MOG::release
 ---------------------
@@ -443,7 +436,9 @@ ocl::MOG2
 -------------
 .. ocv:class:: ocl::MOG2 : public ocl::BackgroundSubtractor
 
-Gaussian Mixture-based Background/Foreground Segmentation Algorithm. ::
+  Gaussian Mixture-based Background/Foreground Segmentation Algorithm.
+
+  The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [MOG2004]_. ::
 
     class CV_EXPORTS MOG2: public cv::ocl::BackgroundSubtractor
     {
@@ -485,45 +480,42 @@ Gaussian Mixture-based Background/Foreground Segmentation Algorithm. ::
         /* hidden */
     };
 
-  The class discriminates between foreground and background pixels by building and maintaining a model of the background. Any pixel which does not fit this model is then deemed to be foreground. The class implements algorithm described in [MOG2004]_.
+  .. ocv:member:: float backgroundRatio
 
-  Here are important members of the class that control the algorithm, which you can set after constructing the class instance:
+      Threshold defining whether the component is significant enough to be included into the background model. ``cf=0.1 => TB=0.9`` is default. For ``alpha=0.001``, it means that the mode should exist for approximately 105 frames before it is considered foreground.
 
-    .. ocv:member:: float backgroundRatio
+  .. ocv:member:: float varThreshold
 
-        Threshold defining whether the component is significant enough to be included into the background model. ``cf=0.1 => TB=0.9`` is default. For ``alpha=0.001``, it means that the mode should exist for approximately 105 frames before it is considered foreground.
+      Threshold for the squared Mahalanobis distance that helps decide when a sample is close to the existing components (corresponds to ``Tg``). If it is not close to any component, a new component is generated. ``3 sigma => Tg=3*3=9`` is default. A smaller ``Tg`` value generates more components. A higher ``Tg`` value may result in a small number of components but they can grow too large.
 
-    .. ocv:member:: float varThreshold
+  .. ocv:member:: float fVarInit
 
-        Threshold for the squared Mahalanobis distance that helps decide when a sample is close to the existing components (corresponds to ``Tg``). If it is not close to any component, a new component is generated. ``3 sigma => Tg=3*3=9`` is default. A smaller ``Tg`` value generates more components. A higher ``Tg`` value may result in a small number of components but they can grow too large.
+      Initial variance for the newly generated components. It affects the speed of adaptation. The parameter value is based on your estimate of the typical standard deviation from the images. OpenCV uses 15 as a reasonable value.
 
-    .. ocv:member:: float fVarInit
+  .. ocv:member:: float fVarMin
 
-        Initial variance for the newly generated components. It affects the speed of adaptation. The parameter value is based on your estimate of the typical standard deviation from the images. OpenCV uses 15 as a reasonable value.
+      Parameter used to further control the variance.
 
-    .. ocv:member:: float fVarMin
+  .. ocv:member:: float fVarMax
 
-        Parameter used to further control the variance.
+      Parameter used to further control the variance.
 
-    .. ocv:member:: float fVarMax
+  .. ocv:member:: float fCT
 
-        Parameter used to further control the variance.
+      Complexity reduction parameter. This parameter defines the number of samples needed to accept to prove the component exists. ``CT=0.05`` is a default value for all the samples. By setting ``CT=0`` you get an algorithm very similar to the standard Stauffer&Grimson algorithm.
 
-    .. ocv:member:: float fCT
+  .. ocv:member:: uchar nShadowDetection
 
-        Complexity reduction parameter. This parameter defines the number of samples needed to accept to prove the component exists. ``CT=0.05`` is a default value for all the samples. By setting ``CT=0`` you get an algorithm very similar to the standard Stauffer&Grimson algorithm.
+      The value for marking shadow pixels in the output foreground mask. Default value is 127.
 
-    .. ocv:member:: uchar nShadowDetection
+  .. ocv:member:: float fTau
 
-        The value for marking shadow pixels in the output foreground mask. Default value is 127.
+      Shadow threshold. The shadow is detected if the pixel is a darker version of the background. ``Tau`` is a threshold defining how much darker the shadow can be. ``Tau= 0.5`` means that if a pixel is more than twice darker then it is not shadow. See [ShadowDetect2003]_.
 
-    .. ocv:member:: float fTau
+  .. ocv:member:: bool bShadowDetection
 
-        Shadow threshold. The shadow is detected if the pixel is a darker version of the background. ``Tau`` is a threshold defining how much darker the shadow can be. ``Tau= 0.5`` means that if a pixel is more than twice darker then it is not shadow. See [ShadowDetect2003]_.
+      Parameter defining whether shadow detection should be enabled.
 
-    .. ocv:member:: bool bShadowDetection
-
-        Parameter defining whether shadow detection should be enabled.
 
 .. seealso:: :ocv:class:`BackgroundSubtractorMOG2`
 
@@ -549,8 +541,6 @@ Updates the background model and returns the foreground mask.
 
     :param fgmask: The output foreground mask as an 8-bit binary image.
 
-    :param stream: Stream for the asynchronous version.
-
 
 ocl::MOG2::getBackgroundImage
 ---------------------------------
@@ -560,11 +550,9 @@ Computes a background image.
 
     :param backgroundImage: The output background image.
 
-    :param stream: Stream for the asynchronous version.
-
 
 ocl::MOG2::release
 ----------------------
 Releases all inner buffer's memory.
 
-.. ocv:function:: void ocl::MOG2::release()
+.. ocv:function:: void ocl::MOG2::release()

modules/ocl/include/opencv2/ocl/ocl.hpp

@@ -308,16 +308,13 @@ namespace cv
             void copyTo( oclMat &m, const oclMat &mask = oclMat()) const;
 
             //! converts oclMatrix to another datatype with optional scalng. See cvConvertScale.
-            //It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
             void convertTo( oclMat &m, int rtype, double alpha = 1, double beta = 0 ) const;
 
             void assignTo( oclMat &m, int type = -1 ) const;
 
             //! sets every oclMatrix element to s
-            //It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
             oclMat& operator = (const Scalar &s);
             //! sets some of the oclMatrix elements to s, according to the mask
-            //It supports 8UC1 8UC4 32SC1 32SC4 32FC1 32FC4
             oclMat& setTo(const Scalar &s, const oclMat &mask = oclMat());
             //! creates alternative oclMatrix header for the same data, with different
             // number of channels and/or different number of rows. see cvReshape.
@@ -556,11 +553,12 @@ namespace cv
         CV_EXPORTS void bilateralFilter(const oclMat& src, oclMat& dst, int d, double sigmaColor, double sigmaSpace, int borderType=BORDER_DEFAULT);
 
         //! Applies an adaptive bilateral filter to the input image
-        //  This is not truly a bilateral filter. Instead of using user provided fixed parameters,
-        //  the function calculates a constant at each window based on local standard deviation,
-        //  and use this constant to do filtering.
+        //  Unlike the usual bilateral filter that uses fixed value for sigmaColor,
+        //  the adaptive version calculates the local variance in he ksize neighborhood
+        //  and use this as sigmaColor, for the value filtering. However, the local standard deviation is
+        //  clamped to the maxSigmaColor.
         //  supports 8UC1, 8UC3
-        CV_EXPORTS void adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, Point anchor = Point(-1, -1), int borderType=BORDER_DEFAULT);
+        CV_EXPORTS void adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, double maxSigmaColor=20.0, Point anchor = Point(-1, -1), int borderType=BORDER_DEFAULT);
 
         //! computes exponent of each matrix element (dst = e**src)
         // supports only CV_32FC1, CV_64FC1 type
@@ -631,6 +629,9 @@ namespace cv
         //! initializes a scaled identity matrix
         CV_EXPORTS void setIdentity(oclMat& src, const Scalar & val = Scalar(1));
 
+        //! fills the output array with repeated copies of the input array
+        CV_EXPORTS void repeat(const oclMat & src, int ny, int nx, oclMat & dst);
+
         //////////////////////////////// Filter Engine ////////////////////////////////
 
         /*!
@@ -898,6 +899,9 @@ namespace cv
 
             CvSeq* oclHaarDetectObjects(oclMat &gimg, CvMemStorage *storage, double scaleFactor,
                                         int minNeighbors, int flags, CvSize minSize = cvSize(0, 0), CvSize maxSize = cvSize(0, 0));
+            void detectMultiScale(oclMat &image, CV_OUT std::vector<cv::Rect>& faces,
+                double scaleFactor = 1.1, int minNeighbors = 3, int flags = 0,
+                Size minSize = Size(), Size maxSize = Size());
         };
 
         class CV_EXPORTS OclCascadeClassifierBuf : public  cv::CascadeClassifier

modules/ocl/perf/perf_arithm.cpp

@@ -1051,3 +1051,40 @@ PERF_TEST_P(AbsFixture, Abs,
     else
         OCL_PERF_ELSE
 }
+
+///////////// Repeat ////////////////////////
+
+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)))
+{
+    const Size_MatType_t params = GetParam();
+    const Size srcSize = get<0>(params);
+    const int type = get<1>(params);
+    const int nx = 3, ny = 2;
+    const Size dstSize(srcSize.width * nx, srcSize.height * ny);
+
+    Mat src(srcSize, type), dst(dstSize, type);
+    declare.in(src, WARMUP_RNG).out(dst);
+
+    if (RUN_OCL_IMPL)
+    {
+        ocl::oclMat oclSrc(src), oclDst(dstSize, type);
+
+        OCL_TEST_CYCLE() cv::ocl::repeat(oclSrc, ny, nx, oclDst);
+
+        oclDst.download(dst);
+
+        SANITY_CHECK(dst);
+    }
+    else if (RUN_PLAIN_IMPL)
+    {
+        TEST_CYCLE() cv::repeat(src, ny, nx, dst);
+
+        SANITY_CHECK(dst);
+    }
+    else
+        OCL_PERF_ELSE
+}

modules/ocl/perf/perf_haar.cpp

@@ -48,44 +48,6 @@
 using namespace perf;
 
 ///////////// Haar ////////////////////////
-namespace cv
-{
-namespace ocl
-{
-
-struct getRect
-{
-    Rect operator()(const CvAvgComp &e) const
-    {
-        return e.rect;
-    }
-};
-
-class CascadeClassifier_GPU : public OclCascadeClassifier
-{
-public:
-    void detectMultiScale(oclMat &image,
-                          CV_OUT std::vector<cv::Rect>& faces,
-                          double scaleFactor = 1.1,
-                          int minNeighbors = 3, int flags = 0,
-                          Size minSize = Size(),
-                          Size maxSize = Size())
-    {
-        (void)maxSize;
-        MemStorage storage(cvCreateMemStorage(0));
-        //CvMat img=image;
-        CvSeq *objs = oclHaarDetectObjects(image, storage, scaleFactor, minNeighbors, flags, minSize);
-        vector<CvAvgComp> vecAvgComp;
-        Seq<CvAvgComp>(objs).copyTo(vecAvgComp);
-        faces.resize(vecAvgComp.size());
-        std::transform(vecAvgComp.begin(), vecAvgComp.end(), faces.begin(), getRect());
-    }
-
-};
-
-}
-}
-
 PERF_TEST(HaarFixture, Haar)
 {
     vector<Rect> faces;
@@ -107,7 +69,7 @@ PERF_TEST(HaarFixture, Haar)
     }
     else if (RUN_OCL_IMPL)
     {
-        ocl::CascadeClassifier_GPU faceCascade;
+        ocl::OclCascadeClassifier faceCascade;
         ocl::oclMat oclImg(img);
 
         ASSERT_TRUE(faceCascade.load(getDataPath("gpu/haarcascade/haarcascade_frontalface_alt.xml")))

modules/ocl/src/arithm.cpp

@@ -1706,3 +1706,21 @@ void cv::ocl::setIdentity(oclMat& src, const Scalar & scalar)
     openCLExecuteKernel(src.clCxt, &arithm_setidentity, "setIdentity", global_threads, local_threads,
                         args, -1, -1, buildOptions.c_str());
 }
+
+//////////////////////////////////////////////////////////////////////////////
+////////////////////////////////// Repeat ////////////////////////////////////
+//////////////////////////////////////////////////////////////////////////////
+
+void cv::ocl::repeat(const oclMat & src, int ny, int nx, oclMat & dst)
+{
+    CV_Assert(nx > 0 && ny > 0);
+    dst.create(src.rows * ny, src.cols * nx, src.type());
+
+    for (int y = 0; y < ny; ++y)
+        for (int x = 0; x < nx; ++x)
+        {
+            Rect roi(x * src.cols, y * src.rows, src.cols, src.rows);
+            oclMat hdr = dst(roi);
+            src.copyTo(hdr);
+        }
+}

modules/ocl/src/cl_operations.cpp

@@ -109,12 +109,15 @@ cl_mem openCLCreateBuffer(Context *ctx, size_t flag , size_t size)
     return buffer;
 }
 
+#define MEMORY_CORRUPTION_GUARD
+#ifdef MEMORY_CORRUPTION_GUARD
 //#define CHECK_MEMORY_CORRUPTION
-#ifdef CHECK_MEMORY_CORRUPTION
-//#define CHECK_MEMORY_CORRUPTION_PRINT_ERROR
+#define CHECK_MEMORY_CORRUPTION_PRINT_ERROR
 #define CHECK_MEMORY_CORRUPTION_RAISE_ERROR
-static const int __memory_corruption_check_bytes = 1024*1024;
+static const int __memory_corruption_guard_bytes = 64*1024;
+#ifdef CHECK_MEMORY_CORRUPTION
 static const int __memory_corruption_check_pattern = 0x14326547; // change pattern for sizeof(int)==8
+#endif
 struct CheckBuffers
 {
     cl_mem mainBuffer;
@@ -128,7 +131,7 @@ struct CheckBuffers
     CheckBuffers(cl_mem _mainBuffer, size_t _size, size_t _widthInBytes, size_t _height)
         : mainBuffer(_mainBuffer), size(_size), widthInBytes(_widthInBytes), height(_height)
     {
-        // notihng
+        // nothing
     }
 };
 static std::map<cl_mem, CheckBuffers> __check_buffers;
@@ -145,32 +148,52 @@ void openCLMallocPitchEx(Context *ctx, void **dev_ptr, size_t *pitch,
 {
     cl_int status;
     size_t size = widthInBytes * height;
-#ifndef CHECK_MEMORY_CORRUPTION
-    *dev_ptr = clCreateBuffer(getClContext(ctx), gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
-                              size, 0, &status);
-    openCLVerifyCall(status);
+    bool useSubBuffers =
+#ifndef MEMORY_CORRUPTION_GUARD
+            false;
 #else
-    size_t allocSize = size + __memory_corruption_check_bytes * 2;
-    cl_mem mainBuffer = clCreateBuffer(getClContext(ctx), gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
-            allocSize, 0, &status);
-    openCLVerifyCall(status);
-    cl_buffer_region r = {__memory_corruption_check_bytes, size};
-    *dev_ptr =  clCreateSubBuffer(mainBuffer,
-            gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
-            CL_BUFFER_CREATE_TYPE_REGION, &r,
-            &status);
-    openCLVerifyCall(status);
-    std::vector<int> tmp(__memory_corruption_check_bytes / sizeof(int),
-            __memory_corruption_check_pattern);
-    CV_Assert(tmp.size() * sizeof(int) == __memory_corruption_check_bytes);
-    openCLVerifyCall(clEnqueueWriteBuffer(getClCommandQueue(ctx),
-            mainBuffer, CL_TRUE, 0, __memory_corruption_check_bytes, &tmp[0],
-            0, NULL, NULL));
-    openCLVerifyCall(clEnqueueWriteBuffer(getClCommandQueue(ctx),
-            mainBuffer, CL_TRUE, __memory_corruption_check_bytes + size, __memory_corruption_check_bytes, &tmp[0],
-            0, NULL, NULL));
-    CheckBuffers data(mainBuffer, size, widthInBytes, height);
-    __check_buffers.insert(std::pair<cl_mem, CheckBuffers>((cl_mem)*dev_ptr, data));
+            true;
+#endif
+    const DeviceInfo& devInfo = ctx->getDeviceInfo();
+    if (useSubBuffers && devInfo.isIntelDevice)
+    {
+        useSubBuffers = false; // TODO FIXIT We observe memory leaks then we working with sub-buffers
+                               // on the CPU device of Intel OpenCL SDK (Linux). We will investigate this later.
+    }
+    if (!useSubBuffers)
+    {
+        *dev_ptr = clCreateBuffer(getClContext(ctx), gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
+                                  size, 0, &status);
+        openCLVerifyCall(status);
+    }
+#ifdef MEMORY_CORRUPTION_GUARD
+    else
+    {
+        size_t allocSize = size + __memory_corruption_guard_bytes * 2;
+        cl_mem mainBuffer = clCreateBuffer(getClContext(ctx), gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
+                allocSize, 0, &status);
+        openCLVerifyCall(status);
+        cl_buffer_region r = {__memory_corruption_guard_bytes, size};
+        *dev_ptr = clCreateSubBuffer(mainBuffer,
+                gDevMemRWValueMap[rw_type]|gDevMemTypeValueMap[mem_type],
+                CL_BUFFER_CREATE_TYPE_REGION, &r,
+                &status);
+        openCLVerifyCall(status);
+#ifdef CHECK_MEMORY_CORRUPTION
+        std::vector<int> tmp(__memory_corruption_guard_bytes / sizeof(int),
+                __memory_corruption_check_pattern);
+        CV_Assert(tmp.size() * sizeof(int) == __memory_corruption_guard_bytes);
+        openCLVerifyCall(clEnqueueWriteBuffer(getClCommandQueue(ctx),
+                mainBuffer, CL_FALSE, 0, __memory_corruption_guard_bytes, &tmp[0],
+                0, NULL, NULL));
+        openCLVerifyCall(clEnqueueWriteBuffer(getClCommandQueue(ctx),
+                mainBuffer, CL_FALSE, __memory_corruption_guard_bytes + size, __memory_corruption_guard_bytes, &tmp[0],
+                0, NULL, NULL));
+        clFinish(getClCommandQueue(ctx));
+#endif
+        CheckBuffers data(mainBuffer, size, widthInBytes, height);
+        __check_buffers.insert(std::pair<cl_mem, CheckBuffers>((cl_mem)*dev_ptr, data));
+    }
 #endif
     *pitch = widthInBytes;
 }
@@ -224,40 +247,48 @@ void openCLCopyBuffer2D(Context *ctx, void *dst, size_t dpitch, int dst_offset,
 
 void openCLFree(void *devPtr)
 {
+    openCLSafeCall(clReleaseMemObject((cl_mem)devPtr));
+#ifdef MEMORY_CORRUPTION_GUARD
 #ifdef CHECK_MEMORY_CORRUPTION
     bool failBefore = false, failAfter = false;
+#endif
     CheckBuffers data;
     std::map<cl_mem, CheckBuffers>::iterator i = __check_buffers.find((cl_mem)devPtr);
     if (i != __check_buffers.end())
     {
         data = i->second;
+#ifdef CHECK_MEMORY_CORRUPTION
         Context* ctx = Context::getContext();
-        std::vector<uchar> checkBefore(__memory_corruption_check_bytes);
-        std::vector<uchar> checkAfter(__memory_corruption_check_bytes);
+        std::vector<uchar> checkBefore(__memory_corruption_guard_bytes);
+        std::vector<uchar> checkAfter(__memory_corruption_guard_bytes);
         openCLVerifyCall(clEnqueueReadBuffer(getClCommandQueue(ctx),
-                data.mainBuffer, CL_TRUE, 0, __memory_corruption_check_bytes, &checkBefore[0],
+                data.mainBuffer, CL_FALSE, 0, __memory_corruption_guard_bytes, &checkBefore[0],
                 0, NULL, NULL));
         openCLVerifyCall(clEnqueueReadBuffer(getClCommandQueue(ctx),
-                data.mainBuffer, CL_TRUE, __memory_corruption_check_bytes + data.size, __memory_corruption_check_bytes, &checkAfter[0],
+                data.mainBuffer, CL_FALSE, __memory_corruption_guard_bytes + data.size, __memory_corruption_guard_bytes, &checkAfter[0],
                 0, NULL, NULL));
+        clFinish(getClCommandQueue(ctx));
 
-        std::vector<int> tmp(__memory_corruption_check_bytes / sizeof(int),
+        std::vector<int> tmp(__memory_corruption_guard_bytes / sizeof(int),
                 __memory_corruption_check_pattern);
 
-        if (memcmp(&checkBefore[0], &tmp[0], __memory_corruption_check_bytes) != 0)
+        if (memcmp(&checkBefore[0], &tmp[0], __memory_corruption_guard_bytes) != 0)
         {
             failBefore = true;
         }
-        if (memcmp(&checkAfter[0], &tmp[0], __memory_corruption_check_bytes) != 0)
+        if (memcmp(&checkAfter[0], &tmp[0], __memory_corruption_guard_bytes) != 0)
         {
             failAfter = true;
         }
+#else
+        // TODO FIXIT Attach clReleaseMemObject call to event completion callback
+        Context* ctx = Context::getContext();
+        clFinish(getClCommandQueue(ctx));
+#endif
         openCLSafeCall(clReleaseMemObject(data.mainBuffer));
         __check_buffers.erase(i);
     }
-#endif
-    openCLSafeCall(clReleaseMemObject((cl_mem)devPtr));
-#ifdef CHECK_MEMORY_CORRUPTION
+#if defined(CHECK_MEMORY_CORRUPTION)
     if (failBefore)
     {
 #ifdef CHECK_MEMORY_CORRUPTION_PRINT_ERROR
@@ -276,7 +307,8 @@ void openCLFree(void *devPtr)
         CV_Error(CV_StsInternal, "Memory corruption detected: after buffer");
 #endif
     }
-#endif
+#endif // CHECK_MEMORY_CORRUPTION
+#endif // MEMORY_CORRUPTION_GUARD
 }
 
 cl_kernel openCLGetKernelFromSource(const Context *ctx, const cv::ocl::ProgramEntry* source, string kernelName)

modules/ocl/src/filtering.cpp

@@ -20,6 +20,7 @@
 //    Zero Lin, Zero.Lin@amd.com
 //    Zhang Ying, zhangying913@gmail.com
 //    Yao Wang, bitwangyaoyao@gmail.com
+//    Harris Gasparakis, harris.gasparakis@amd.com
 //
 // Redistribution and use in source and binary forms, with or without modification,
 // are permitted provided that the following conditions are met:
@@ -1407,7 +1408,7 @@ void cv::ocl::GaussianBlur(const oclMat &src, oclMat &dst, Size ksize, double si
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 // Adaptive Bilateral Filter
 
-void cv::ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, Point anchor, int borderType)
+void cv::ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize, double sigmaSpace, double maxSigmaColor, Point anchor, int borderType)
 {
     CV_Assert((ksize.width & 1) && (ksize.height & 1));  // ksize must be odd
     CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC3);  // source must be 8bit RGB image
@@ -1418,10 +1419,24 @@ void cv::ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize
     int idx = 0;
     int w = ksize.width / 2;
     int h = ksize.height / 2;
-    for(int y=-h; y<=h; y++)
-        for(int x=-w; x<=w; x++)
+
+    int ABF_GAUSSIAN_ocl = 1;
+
+    if(ABF_GAUSSIAN_ocl)
     {
-        lut.at<float>(idx++) = sigma2 / (sigma2 + x * x + y * y);
+        for(int y=-h; y<=h; y++)
+            for(int x=-w; x<=w; x++)
+        {
+            lut.at<float>(idx++) = expf( (float)(-0.5 * (x * x + y * y)/sigma2));
+        }
+    }
+    else
+    {
+        for(int y=-h; y<=h; y++)
+            for(int x=-w; x<=w; x++)
+        {
+            lut.at<float>(idx++) = (float) (sigma2 / (sigma2 + x * x + y * y));
+        }
     }
 
     oclMat dlut(lut);
@@ -1429,7 +1444,7 @@ void cv::ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize
     int cn = src.oclchannels();
 
     normalizeAnchor(anchor, ksize);
-    const static String kernelName = "edgeEnhancingFilter";
+    const static String kernelName = "adaptiveBilateralFilter";
 
     dst.create(src.size(), src.type());
 
@@ -1478,9 +1493,10 @@ void cv::ocl::adaptiveBilateralFilter(const oclMat& src, oclMat& dst, Size ksize
 
     //LDATATYPESIZE is sizeof local data store. This is to exemplify effect of LDS on kernel performance
     sprintf(build_options,
-        "-D VAR_PER_CHANNEL=1 -D CALCVAR=1 -D FIXED_WEIGHT=0 -D EXTRA=%d"
+        "-D VAR_PER_CHANNEL=1 -D CALCVAR=1 -D FIXED_WEIGHT=0 -D EXTRA=%d -D MAX_VAR_VAL=%f -D ABF_GAUSSIAN=%d"
         " -D THREADS=%d -D anX=%d -D anY=%d -D ksX=%d -D ksY=%d -D %s",
-        static_cast<int>(EXTRA), static_cast<int>(blockSizeX), anchor.x, anchor.y, ksize.width, ksize.height, btype);
+        static_cast<int>(EXTRA), static_cast<float>(maxSigmaColor*maxSigmaColor), static_cast<int>(ABF_GAUSSIAN_ocl),
+        static_cast<int>(blockSizeX), anchor.x, anchor.y, ksize.width, ksize.height, btype);
 
     std::vector<pair<size_t , const void *> > args;
     args.push_back(std::make_pair(sizeof(cl_mem), &src.data));

modules/ocl/src/haar.cpp

@@ -1186,6 +1186,28 @@ CvSeq *cv::ocl::OclCascadeClassifier::oclHaarDetectObjects( oclMat &gimg, CvMemS
     return result_seq;
 }
 
+
+struct getRect
+{
+    Rect operator()(const CvAvgComp &e) const
+    {
+        return e.rect;
+    }
+};
+
+void cv::ocl::OclCascadeClassifier::detectMultiScale(oclMat &gimg, CV_OUT std::vector<cv::Rect>& faces,
+                                                        double scaleFactor, int minNeighbors, int flags,
+                                                        Size minSize, Size maxSize)
+{
+    CvSeq* _objects;
+    MemStorage storage(cvCreateMemStorage(0));
+    _objects = oclHaarDetectObjects(gimg, storage, scaleFactor, minNeighbors, flags, minSize, maxSize);
+    vector<CvAvgComp> vecAvgComp;
+    Seq<CvAvgComp>(_objects).copyTo(vecAvgComp);
+    faces.resize(vecAvgComp.size());
+    std::transform(vecAvgComp.begin(), vecAvgComp.end(), faces.begin(), getRect());
+}
+
 struct OclBuffers
 {
     cl_mem stagebuffer;
@@ -1197,13 +1219,6 @@ struct OclBuffers
     cl_mem newnodebuffer;
 };
 
-struct getRect
-{
-    Rect operator()(const CvAvgComp &e) const
-    {
-        return e.rect;
-    }
-};
 
 void cv::ocl::OclCascadeClassifierBuf::detectMultiScale(oclMat &gimg, CV_OUT std::vector<cv::Rect>& faces,
                                                         double scaleFactor, int minNeighbors, int flags,

modules/ocl/src/opencl/bgfg_mog.cl

@@ -48,22 +48,22 @@
 #define T_MEAN_VAR float
 #define CONVERT_TYPE convert_uchar_sat
 #define F_ZERO (0.0f)
-float cvt(uchar val)
+inline float cvt(uchar val)
 {
     return val;
 }
 
-float sqr(float val)
+inline float sqr(float val)
 {
     return val * val;
 }
 
-float sum(float val)
+inline float sum(float val)
 {
     return val;
 }
 
-float clamp1(float var, float learningRate, float diff, float minVar)
+static float clamp1(float var, float learningRate, float diff, float minVar)
 {
     return fmax(var + learningRate * (diff * diff - var), minVar);
 }
@@ -72,7 +72,7 @@ float clamp1(float var, float learningRate, float diff, float minVar)
 #define T_MEAN_VAR float4
 #define CONVERT_TYPE convert_uchar4_sat
 #define F_ZERO (0.0f, 0.0f, 0.0f, 0.0f)
-float4 cvt(const uchar4 val)
+inline float4 cvt(const uchar4 val)
 {
     float4 result;
     result.x = val.x;
@@ -83,17 +83,17 @@ float4 cvt(const uchar4 val)
     return result;
 }
 
-float sqr(const float4 val)
+inline float sqr(const float4 val)
 {
     return val.x * val.x + val.y * val.y + val.z * val.z;
 }
 
-float sum(const float4 val)
+inline float sum(const float4 val)
 {
     return (val.x + val.y + val.z);
 }
 
-float4 clamp1(const float4 var, float learningRate, const float4 diff, float minVar)
+static float4 clamp1(const float4 var, float learningRate, const float4 diff, float minVar)
 {
     float4 result;
     result.x = fmax(var.x + learningRate * (diff.x * diff.x - var.x), minVar);
@@ -116,14 +116,14 @@ typedef struct
     uchar c_shadowVal;
 }con_srtuct_t;
 
-void swap(__global float* ptr, int x, int y, int k, int rows, int ptr_step)
+static void swap(__global float* ptr, int x, int y, int k, int rows, int ptr_step)
 {
     float val = ptr[(k * rows + y) * ptr_step + x];
     ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x];
     ptr[((k + 1) * rows + y) * ptr_step + x] = val;
 }
 
-void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_step)
+static void swap4(__global float4* ptr, int x, int y, int k, int rows, int ptr_step)
 {
     float4 val = ptr[(k * rows + y) * ptr_step + x];
     ptr[(k * rows + y) * ptr_step + x] = ptr[((k + 1) * rows + y) * ptr_step + x];
@@ -412,7 +412,7 @@ __kernel void mog2_kernel(__global T_FRAME * frame, __global int* fgmask, __glob
 
             if (_weight < -prune)
             {
-                _weight = 0.0;
+                _weight = 0.0f;
                 nmodes--;
             }
 

modules/ocl/src/opencl/filtering_adaptive_bilateral.cl

@@ -85,7 +85,7 @@
 #endif
 
 __kernel void
-edgeEnhancingFilter_C4_D0(
+adaptiveBilateralFilter_C4_D0(
     __global const uchar4 * restrict src,
     __global uchar4 *dst,
     float alpha,
@@ -173,14 +173,14 @@ edgeEnhancingFilter_C4_D0(
         //find variance of all data
         int startLMj;
         int endLMj ;
-#if CALCVAR
         // Top row: don't sum the very last element
         for(int extraCnt = 0; extraCnt <=EXTRA; extraCnt++)
         {
+#if CALCVAR
             startLMj = extraCnt;
             endLMj =  ksY+extraCnt-1;
-            sumVal =0;
-            sumValSqr=0;
+            sumVal = (int4)0;
+            sumValSqr= (int4)0;
             for(int j = startLMj; j < endLMj; j++)
                 for(int i=-anX; i<=anX; i++)
                 {
@@ -190,9 +190,10 @@ edgeEnhancingFilter_C4_D0(
                     sumValSqr += mul24(currVal, currVal);
                 }
 
-            var[extraCnt] = convert_float4( ( (sumValSqr * howManyAll)- mul24(sumVal , sumVal) ) ) /  ( (float)(howManyAll*howManyAll) ) ;
+            var[extraCnt] = clamp( convert_float4( ( (sumValSqr * howManyAll)- mul24(sumVal , sumVal) ) ) /  ( (float)(howManyAll*howManyAll) ), (float4)(0.1f, 0.1f, 0.1f, 0.1f), (float4)(MAX_VAR_VAL, MAX_VAR_VAL, MAX_VAR_VAL, MAX_VAR_VAL)) ;
+
 #else
-        var[extraCnt] = (float4)(900.0, 900.0, 900.0, 0.0);
+            var[extraCnt] = (float4)(MAX_VAR_VAL, MAX_VAR_VAL, MAX_VAR_VAL, MAX_VAR_VAL);
 #endif
         }
 
@@ -221,32 +222,48 @@ edgeEnhancingFilter_C4_D0(
 #else
                     weight = 1.0f;
 #endif
-#else
+#else // !FIXED_WEIGHT
                     currVal = convert_int4(data[j][col+anX+i]);
                     currWRTCenter = currVal-currValCenter;
 
+#if ABF_GAUSSIAN
+
+#if VAR_PER_CHANNEL
+                    weight = exp( (float4)(-0.5f, -0.5f, -0.5f, -0.5f) * convert_float4(currWRTCenter * currWRTCenter) / var[extraCnt] )*
+                        (float4)(lut[lut_j*lut_step+anX+i]);
+#else
+                    weight = exp( -0.5f * (mul24(currWRTCenter.x, currWRTCenter.x) + mul24(currWRTCenter.y, currWRTCenter.y) +
+                        mul24(currWRTCenter.z, currWRTCenter.z) ) / (var[extraCnt].x+var[extraCnt].y+var[extraCnt].z) ) * lut[lut_j*lut_step+anX+i];
+#endif
+
+#else // !ABF_GAUSSIAN
+
 #if VAR_PER_CHANNEL
                     weight = var[extraCnt] / (var[extraCnt] + convert_float4(currWRTCenter * currWRTCenter)) *
                         (float4)(lut[lut_j*lut_step+anX+i]);
 #else
-                    weight = 1.0f/(1.0f+( mul24(currWRTCenter.x, currWRTCenter.x) + mul24(currWRTCenter.y, currWRTCenter.y) +
-                        mul24(currWRTCenter.z, currWRTCenter.z))/(var.x+var.y+var.z));
-#endif
+                    weight = ((float)lut[lut_j*lut_step+anX+i]) /(1.0f+( mul24(currWRTCenter.x, currWRTCenter.x) + mul24(currWRTCenter.y, currWRTCenter.y) +
+                        mul24(currWRTCenter.z, currWRTCenter.z))/(var[extraCnt].x+var[extraCnt].y+var[extraCnt].z));
 #endif
+
+#endif //ABF_GAUSSIAN
+
+
+
+#endif  // FIXED_WEIGHT
+
                     tmp_sum[extraCnt] += convert_float4(data[j][col+anX+i]) * weight;
                     totalWeight += weight;
                 }
             }
 
-            tmp_sum[extraCnt] /= totalWeight;
-
             if(posX >= 0 && posX < dst_cols && (posY+extraCnt) >= 0 && (posY+extraCnt) < dst_rows)
-                dst[(dst_startY+extraCnt) * (dst_step>>2)+ dst_startX + col] = convert_uchar4(tmp_sum[extraCnt]);
+                dst[(dst_startY+extraCnt) * (dst_step>>2)+ dst_startX + col] = convert_uchar4_rtz( (tmp_sum[extraCnt] / (float4)totalWeight) + (float4)0.5f);
 
 #if VAR_PER_CHANNEL
             totalWeight = (float4)(0,0,0,0);
 #else
-            totalWeight = 0;
+            totalWeight = 0.0f;
 #endif
         }
     }
@@ -254,7 +271,7 @@ edgeEnhancingFilter_C4_D0(
 
 
 __kernel void
-edgeEnhancingFilter_C1_D0(
+adaptiveBilateralFilter_C1_D0(
     __global const uchar * restrict src,
     __global uchar *dst,
     float alpha,
@@ -343,10 +360,11 @@ edgeEnhancingFilter_C1_D0(
         //find variance of all data
         int startLMj;
         int endLMj;
-#if CALCVAR
+
         // Top row: don't sum the very last element
         for(int extraCnt=0; extraCnt<=EXTRA; extraCnt++)
         {
+#if CALCVAR
             startLMj = extraCnt;
             endLMj =  ksY+extraCnt-1;
             sumVal = 0;
@@ -361,9 +379,9 @@ edgeEnhancingFilter_C1_D0(
                     sumValSqr += mul24(currVal, currVal);
                 }
             }
-            var[extraCnt] = (float)( ( (sumValSqr * howManyAll)- mul24(sumVal , sumVal) ) ) /  ( (float)(howManyAll*howManyAll) ) ;
+            var[extraCnt] =  clamp((float)( ( (sumValSqr * howManyAll)- mul24(sumVal , sumVal) ) ) /  ( (float)(howManyAll*howManyAll) ) , 0.1f, (float)(MAX_VAR_VAL) );
 #else
-        var[extraCnt] = (float)(900.0);
+            var[extraCnt] = (float)(MAX_VAR_VAL);
 #endif
         }
 
@@ -389,19 +407,20 @@ edgeEnhancingFilter_C1_D0(
                     currVal	= (int)(data[j][col+anX+i])	;
                     currWRTCenter = currVal-currValCenter;
 
+#if ABF_GAUSSIAN
+                    weight = exp( -0.5f * (float)mul24(currWRTCenter,currWRTCenter)/var[extraCnt]) * lut[lut_j*lut_step+anX+i] ;
+#else
                     weight = var[extraCnt] / (var[extraCnt] + (float)mul24(currWRTCenter,currWRTCenter)) * lut[lut_j*lut_step+anX+i] ;
+#endif
 #endif
                     tmp_sum[extraCnt] += (float)(data[j][col+anX+i] * weight);
                     totalWeight += weight;
                 }
             }
 
-            tmp_sum[extraCnt] /= totalWeight;
-
-
             if(posX >= 0 && posX < dst_cols && (posY+extraCnt) >= 0 && (posY+extraCnt) < dst_rows)
             {
-                dst[(dst_startY+extraCnt) * (dst_step)+ dst_startX + col] = (uchar)(tmp_sum[extraCnt]);
+                dst[(dst_startY+extraCnt) * (dst_step)+ dst_startX + col] = convert_uchar_rtz(tmp_sum[extraCnt]/totalWeight+0.5f);
             }
 
             totalWeight = 0;

modules/ocl/src/opencl/haarobjectdetect.cl

@@ -292,7 +292,6 @@ __kernel void __attribute__((reqd_work_group_size(8,8,1)))gpuRunHaarClassifierCa
     for(int scalei = 0; scalei <loopcount; scalei++)
     {
         int4 scaleinfo1= info[scalei];
-        int width = (scaleinfo1.x & 0xffff0000) >> 16;
         int height = scaleinfo1.x & 0xffff;
         int grpnumperline =(scaleinfo1.y & 0xffff0000) >> 16;
         int totalgrp = scaleinfo1.y & 0xffff;

modules/ocl/src/opencl/haarobjectdetect_scaled2.cl

@@ -136,8 +136,6 @@ __kernel void gpuRunHaarClassifierCascade_scaled2(
     {
         int4 scaleinfo1;
         scaleinfo1 = info[scalei];
-        int width = (scaleinfo1.x & 0xffff0000) >> 16;
-        int height = scaleinfo1.x & 0xffff;
         int grpnumperline = (scaleinfo1.y & 0xffff0000) >> 16;
         int totalgrp = scaleinfo1.y & 0xffff;
         float factor = as_float(scaleinfo1.w);

modules/ocl/src/opencl/imgproc_calcHarris.cl

@@ -119,17 +119,17 @@ __kernel void calcHarris(__global const float *Dx, __global const float *Dy, __g
     __local float temp[6][THREADS];
 
 #ifdef BORDER_CONSTANT
-    bool dx_con,dy_con;
-    float dx_s, dy_s;
     for (int i=0; i < ksY+1; i++)
     {
-        dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows;
-        dx_s = Dx[(dx_startY+i)*(dx_step>>2)+(dx_startX+col)];
-        dx_data[i] = dx_con ? dx_s : 0.0f;
+        bool dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows;
+        int indexDx = (dx_startY+i)*(dx_step>>2)+(dx_startX+col);
+        float dx_s = dx_con ? Dx[indexDx] : 0.0f;
+        dx_data[i] = dx_s;
 
-        dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows;
-        dy_s = Dy[(dy_startY+i)*(dy_step>>2)+(dy_startX+col)];
-        dy_data[i] = dy_con ? dy_s : 0.0f;
+        bool dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows;
+        int indexDy = (dy_startY+i)*(dy_step>>2)+(dy_startX+col);
+        float dy_s = dy_con ? Dy[indexDy] : 0.0f;
+        dy_data[i] = dy_s;
 
         data[0][i] = dx_data[i] * dx_data[i];
         data[1][i] = dx_data[i] * dy_data[i];

modules/ocl/src/opencl/imgproc_calcMinEigenVal.cl

@@ -118,16 +118,18 @@ __kernel void calcMinEigenVal(__global const float *Dx,__global const float *Dy,
     __local float temp[6][THREADS];
 
 #ifdef BORDER_CONSTANT
-    bool dx_con, dy_con;
-    float dx_s, dy_s;
     for (int i=0; i < ksY+1; i++)
     {
-        dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows;
-        dx_s = Dx[(dx_startY+i)*(dx_step>>2)+(dx_startX+col)];
-        dx_data[i] = dx_con ? dx_s : 0.0f;
-        dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows;
-        dy_s = Dy[(dy_startY+i)*(dy_step>>2)+(dy_startX+col)];
-        dy_data[i] = dy_con ? dy_s : 0.0f;
+        bool dx_con = dx_startX+col >= 0 && dx_startX+col < dx_whole_cols && dx_startY+i >= 0 && dx_startY+i < dx_whole_rows;
+        int indexDx = (dx_startY+i)*(dx_step>>2)+(dx_startX+col);
+        float dx_s = dx_con ? Dx[indexDx] : 0.0f;
+        dx_data[i] = dx_s;
+
+        bool dy_con = dy_startX+col >= 0 && dy_startX+col < dy_whole_cols && dy_startY+i >= 0 && dy_startY+i < dy_whole_rows;
+        int indexDy = (dy_startY+i)*(dy_step>>2)+(dy_startX+col);
+        float dy_s = dy_con ? Dy[indexDy] : 0.0f;
+        dy_data[i] = dy_s;
+
         data[0][i] = dx_data[i] * dx_data[i];
         data[1][i] = dx_data[i] * dy_data[i];
         data[2][i] = dy_data[i] * dy_data[i];

modules/ocl/src/opencl/tvl1flow.cl

@@ -69,23 +69,16 @@ __global float* dx, __global float* dy, int dx_step)
 
 }
 
-float bicubicCoeff(float x_)
+static float bicubicCoeff(float x_)
 {
 
     float x = fabs(x_);
     if (x <= 1.0f)
-    {
         return x * x * (1.5f * x - 2.5f) + 1.0f;
-    }
     else if (x < 2.0f)
-    {
         return x * (x * (-0.5f * x + 2.5f) - 4.0f) + 2.0f;
-    }
     else
-    {
         return 0.0f;
-    }
-
 }
 
 __kernel void warpBackwardKernel(__global const float* I0, int I0_step, int I0_col, int I0_row,
@@ -170,12 +163,10 @@ __kernel void warpBackwardKernel(__global const float* I0, int I0_step, int I0_c
 
 }
 
-float readImage(__global const float *image,  const int x,  const int y,  const int rows,  const int cols, const int elemCntPerRow)
+static float readImage(__global const float *image,  const int x,  const int y,  const int rows,  const int cols, const int elemCntPerRow)
 {
     int i0 = clamp(x, 0, cols - 1);
     int j0 = clamp(y, 0, rows - 1);
-    int i1 = clamp(x + 1, 0, cols - 1);
-    int j1 = clamp(y + 1, 0, rows - 1);
 
     return image[j0 * elemCntPerRow + i0];
 }
@@ -303,7 +294,7 @@ __kernel void estimateDualVariablesKernel(__global const float* u1, int u1_col,
 
 }
 
-float divergence(__global const float* v1, __global const float* v2, int y, int x, int v1_step, int v2_step)
+static float divergence(__global const float* v1, __global const float* v2, int y, int x, int v1_step, int v2_step)
 {
 
     if (x > 0 && y > 0)
@@ -407,5 +398,4 @@ __kernel void estimateUKernel(__global const float* I1wx, int I1wx_col, int I1wx
             error[y * I1wx_step + x] = n1 + n2;
         }
     }
-
 }

modules/ocl/src/svm.cpp

@@ -686,9 +686,6 @@ float CvSVM_OCL::predict(const CvMat* samples, CV_OUT CvMat* results) const
     }
 
 #else
-    // TODO fix it
-    CV_Error(CV_StsNotImplemented, "This part of code contains mistakes. Install AMD BLAS in order to get a correct result or use CPU version of SVM");
-
     double degree1 = 0.0;
     if (params.kernel_type == CvSVM::POLY)
         degree1 = params.degree;
@@ -813,9 +810,6 @@ bool CvSVMSolver_ocl::solve_generic( CvSVMSolutionInfo& si )
     }
 
 #else
-    // TODO fix it
-    CV_Error(CV_StsNotImplemented, "This part of code contains mistakes. Install AMD BLAS in order to get a correct result or use CPU version of SVM");
-
     double degree1 = 0.0;
     if(params->kernel_type == CvSVM::POLY)
         degree1 = params->degree;
@@ -1000,13 +994,15 @@ void CvSVMKernel_ocl::calc( int vcount, const int row_idx, Qfloat* results, Mat&
     //int j;
     (this->*calc_func_ocl)( vcount, row_idx, results, src);
 
-// FIXIT #if defined HAVE_CLAMDBLAS
+#if !defined(HAVE_CLAMDBLAS)
+    // nothing
+#else
     const Qfloat max_val = (Qfloat)(FLT_MAX * 1e-3);
     int j;
     for( j = 0; j < vcount; j++ )
         if( results[j] > max_val )
             results[j] = max_val;
-// FIXIT #endif
+#endif
 }
 
 bool CvSVMKernel_ocl::create( const CvSVMParams* _params, Calc_ocl _calc_func, Calc _calc_func1 )
@@ -1078,12 +1074,13 @@ void CvSVMKernel_ocl::calc_poly( int vcount, const int row_idx, Qfloat* results,
 {
     calc_non_rbf_base( vcount, row_idx, results, src);
 
-//FIXIT #if defined HAVE_CLAMDBLAS
-
+#if !defined(HAVE_CLAMDBLAS)
+    // nothing
+#else
     CvMat R = cvMat( 1, vcount, QFLOAT_TYPE, results );
     if( vcount > 0 )
         cvPow( &R, &R, params->degree );
-//FIXIT #endif
+#endif
 }
 
 
@@ -1091,7 +1088,9 @@ void CvSVMKernel_ocl::calc_sigmoid( int vcount, const int row_idx, Qfloat* resul
 {
     calc_non_rbf_base( vcount, row_idx, results, src);
     // TODO: speedup this
-//FIXIT #if defined HAVE_CLAMDBLAS
+#if !defined(HAVE_CLAMDBLAS)
+    // nothing
+#else
     for(int j = 0; j < vcount; j++ )
     {
         Qfloat t = results[j];
@@ -1101,7 +1100,7 @@ void CvSVMKernel_ocl::calc_sigmoid( int vcount, const int row_idx, Qfloat* resul
         else
             results[j] = (Qfloat)((e - 1.) / (e + 1.));
     }
-//FIXIT #endif
+#endif
 }
 
 CvSVM_OCL::CvSVM_OCL()

modules/ocl/test/test_arithm.cpp

@@ -192,13 +192,13 @@ PARAM_TEST_CASE(ArithmTestBase, MatDepth, Channels, bool)
         use_roi = GET_PARAM(2);
     }
 
-    void random_roi()
+    virtual void random_roi()
     {
         const int type = CV_MAKE_TYPE(depth, cn);
 
         Size roiSize = randomSize(1, MAX_VALUE);
-        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
-        randomSubMat(src1, src1_roi, roiSize, srcBorder, type, 2, 11);
+        Border src1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src1, src1_roi, roiSize, src1Border, type, 2, 11);
 
         Border src2Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
         randomSubMat(src2, src2_roi, roiSize, src2Border, type, -1540, 1740);
@@ -214,7 +214,7 @@ PARAM_TEST_CASE(ArithmTestBase, MatDepth, Channels, bool)
         cv::threshold(mask, mask, 0.5, 255., CV_8UC1);
 
 
-        generateOclMat(gsrc1_whole, gsrc1_roi, src1, roiSize, srcBorder);
+        generateOclMat(gsrc1_whole, gsrc1_roi, src1, roiSize, src1Border);
         generateOclMat(gsrc2_whole, gsrc2_roi, src2, roiSize, src2Border);
         generateOclMat(gdst1_whole, gdst1_roi, dst1, roiSize, dst1Border);
         generateOclMat(gdst2_whole, gdst2_roi, dst2, roiSize, dst2Border);
@@ -1522,6 +1522,48 @@ OCL_TEST_P(Norm, NORM_L2)
         }
 }
 
+//// Repeat
+
+struct RepeatTestCase :
+        public ArithmTestBase
+{
+    int nx, ny;
+
+    virtual void random_roi()
+    {
+        const int type = CV_MAKE_TYPE(depth, cn);
+
+        nx = randomInt(1, 4);
+        ny = randomInt(1, 4);
+
+        Size srcRoiSize = randomSize(1, MAX_VALUE);
+        Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(src1, src1_roi, srcRoiSize, srcBorder, type, 2, 11);
+
+        Size dstRoiSize(srcRoiSize.width * nx, srcRoiSize.height * ny);
+        Border dst1Border = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(dst1, dst1_roi, dstRoiSize, dst1Border, type, 5, 16);
+
+        generateOclMat(gsrc1_whole, gsrc1_roi, src1, srcRoiSize, srcBorder);
+        generateOclMat(gdst1_whole, gdst1_roi, dst1, dstRoiSize, dst1Border);
+    }
+};
+
+typedef RepeatTestCase Repeat;
+
+OCL_TEST_P(Repeat, Mat)
+{
+    for (int i = 0; i < LOOP_TIMES; ++i)
+    {
+        random_roi();
+
+        cv::repeat(src1_roi, ny, nx, dst1_roi);
+        cv::ocl::repeat(gsrc1_roi, ny, nx, gdst1_roi);
+
+        Near();
+    }
+}
+
 //////////////////////////////////////// Instantiation /////////////////////////////////////////
 
 INSTANTIATE_TEST_CASE_P(Arithm, Lut, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool(), Bool()));
@@ -1557,5 +1599,6 @@ INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, Combine(Values(CV_8U, CV_8S, CV_16U
 INSTANTIATE_TEST_CASE_P(Arithm, SetIdentity, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
 INSTANTIATE_TEST_CASE_P(Arithm, MeanStdDev, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
 INSTANTIATE_TEST_CASE_P(Arithm, Norm, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
+INSTANTIATE_TEST_CASE_P(Arithm, Repeat, Combine(Values(CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F), Values(1, 2, 3, 4), Bool()));
 
 #endif // HAVE_OPENCL

modules/ocl/test/test_filters.cpp

@@ -338,8 +338,8 @@ OCL_TEST_P(AdaptiveBilateral, Mat)
     {
         random_roi();
 
-        adaptiveBilateralFilter(src_roi, dst_roi, kernelSize, 5, Point(-1, -1), borderType); // TODO anchor
-        ocl::adaptiveBilateralFilter(gsrc_roi, gdst_roi, kernelSize, 5, Point(-1, -1), borderType);
+        adaptiveBilateralFilter(src_roi, dst_roi, kernelSize, 5, 1, Point(-1, -1), borderType); // TODO anchor
+        ocl::adaptiveBilateralFilter(gsrc_roi, gdst_roi, kernelSize, 5, 1, Point(-1, -1), borderType);
 
         Near();
     }

modules/ocl/test/test_imgproc.cpp

@@ -93,14 +93,22 @@ PARAM_TEST_CASE(ImgprocTestBase, MatType,
         generateOclMat(gdst_whole, gdst_roi, dst_whole, roiSize, dstBorder);
     }
 
-    void Near(double threshold = 0.0)
+    void Near(double threshold = 0.0, bool relative = false)
     {
-        Mat whole, roi;
+        Mat roi, whole;
         gdst_whole.download(whole);
         gdst_roi.download(roi);
 
-        EXPECT_MAT_NEAR(dst_whole, whole, threshold);
-        EXPECT_MAT_NEAR(dst_roi, roi, threshold);
+        if (relative)
+        {
+            EXPECT_MAT_NEAR_RELATIVE(dst_whole, whole, threshold);
+            EXPECT_MAT_NEAR_RELATIVE(dst_roi, roi, threshold);
+        }
+        else
+        {
+            EXPECT_MAT_NEAR(dst_whole, whole, threshold);
+            EXPECT_MAT_NEAR(dst_roi, roi, threshold);
+        }
     }
 };
 
@@ -134,18 +142,23 @@ PARAM_TEST_CASE(CopyMakeBorder, MatDepth, // depth
 
     void random_roi()
     {
+        border = randomBorder(0, MAX_VALUE << 2);
+        val = randomScalar(-MAX_VALUE, MAX_VALUE);
+
         Size roiSize = randomSize(1, MAX_VALUE);
         Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
-        randomSubMat(src, src_roi, roiSize, srcBorder, type, 5, 256);
+        randomSubMat(src, src_roi, roiSize, srcBorder, type, -MAX_VALUE, MAX_VALUE);
 
         Border dstBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
-        randomSubMat(dst_whole, dst_roi, roiSize, dstBorder, type, 5, 16);
+        dstBorder.top += border.top;
+        dstBorder.lef += border.lef;
+        dstBorder.rig += border.rig;
+        dstBorder.bot += border.bot;
+
+        randomSubMat(dst_whole, dst_roi, roiSize, dstBorder, type, -MAX_VALUE, MAX_VALUE);
 
         generateOclMat(gsrc_whole, gsrc_roi, src, roiSize, srcBorder);
         generateOclMat(gdst_whole, gdst_roi, dst_whole, roiSize, dstBorder);
-
-        border = randomBorder(0, MAX_VALUE << 2);
-        val = randomScalar(-MAX_VALUE, MAX_VALUE);
     }
 
     void Near(double threshold = 0.0)
@@ -199,11 +212,19 @@ struct CornerTestBase :
         Mat image = readImageType("gpu/stereobm/aloe-L.png", type);
         ASSERT_FALSE(image.empty());
 
+        bool isFP = CV_MAT_DEPTH(type) >= CV_32F;
+        float val = 255.0f;
+        if (isFP)
+        {
+            image.convertTo(image, -1, 1.0 / 255);
+            val /= 255.0f;
+        }
+
         Size roiSize = image.size();
         Border srcBorder = randomBorder(0, useRoi ? MAX_VALUE : 0);
 
         Size wholeSize = Size(roiSize.width + srcBorder.lef + srcBorder.rig, roiSize.height + srcBorder.top + srcBorder.bot);
-        src = randomMat(wholeSize, type, -255, 255, false);
+        src = randomMat(wholeSize, type, -val, val, false);
         src_roi = src(Rect(srcBorder.lef, srcBorder.top, roiSize.width, roiSize.height));
         image.copyTo(src_roi);
 
@@ -228,7 +249,7 @@ OCL_TEST_P(CornerMinEigenVal, Mat)
         cornerMinEigenVal(src_roi, dst_roi, blockSize, apertureSize, borderType);
         ocl::cornerMinEigenVal(gsrc_roi, gdst_roi, blockSize, apertureSize, borderType);
 
-        Near(0.02);
+        Near(1e-5, true);
     }
 }
 
@@ -248,7 +269,7 @@ OCL_TEST_P(CornerHarris, Mat)
         cornerHarris(src_roi, dst_roi, blockSize, apertureSize, k, borderType);
         ocl::cornerHarris(gsrc_roi, gdst_roi, blockSize, apertureSize, k, borderType);
 
-        Near(0.02);
+        Near(1e-5, true);
     }
 }
 
@@ -514,7 +535,7 @@ INSTANTIATE_TEST_CASE_P(Imgproc, CornerMinEigenVal, Combine(
                             Bool()));
 
 INSTANTIATE_TEST_CASE_P(Imgproc, CornerHarris, Combine(
-                            Values((MatType)CV_8UC1), // TODO does not work properly with CV_32FC1
+                            Values((MatType)CV_8UC1, CV_32FC1),
                             Values(3, 5),
                             Values( (int)BORDER_CONSTANT, (int)BORDER_REPLICATE, (int)BORDER_REFLECT, (int)BORDER_REFLECT_101),
                             Bool()));
@@ -559,14 +580,11 @@ INSTANTIATE_TEST_CASE_P(Imgproc, ColumnSum, Combine(
                             Bool()));
 
 INSTANTIATE_TEST_CASE_P(ImgprocTestBase, CopyMakeBorder, Combine(
-                            testing::Range((MatDepth)CV_8U, (MatDepth)CV_USRTYPE1),
-                            testing::Values((Channels)1, (Channels)4),
+                            testing::Values((MatDepth)CV_8U, (MatDepth)CV_16S, (MatDepth)CV_32S, (MatDepth)CV_32F),
+                            testing::Values(Channels(1), Channels(3), (Channels)4),
                             Bool(), // border isolated or not
-                            Values((Border)BORDER_CONSTANT,
-                                   (Border)BORDER_REPLICATE,
-                                   (Border)BORDER_REFLECT,
-                                   (Border)BORDER_WRAP,
-                                   (Border)BORDER_REFLECT_101),
+                            Values((Border)BORDER_REPLICATE, (Border)BORDER_REFLECT,
+                                   (Border)BORDER_WRAP, (Border)BORDER_REFLECT_101),
                             Bool()));
 
 #endif // HAVE_OPENCL

modules/ocl/test/test_ml.cpp

@@ -126,8 +126,6 @@ OCL_TEST_P(KNN, Accuracy)
 INSTANTIATE_TEST_CASE_P(OCL_ML, KNN, Combine(Values(6, 5), Values(Size(200, 400), Size(300, 600)),
     Values(4, 3), Values(false, true)));
 
-#ifdef HAVE_CLAMDBLAS // TODO does not work non-blas version of SVM
-
 ////////////////////////////////SVM/////////////////////////////////////////////////
 
 PARAM_TEST_CASE(SVM_OCL, int, int, int)
@@ -308,6 +306,4 @@ INSTANTIATE_TEST_CASE_P(OCL_ML, SVM_OCL, testing::Combine(
                             Values(2, 3, 4)
                         ));
 
-#endif // HAVE_CLAMDBLAS
-
 #endif // HAVE_OPENCL

modules/ocl/test/test_objdetect.cpp

@@ -218,14 +218,9 @@ PARAM_TEST_CASE(Haar, int, CascadeName)
 
 OCL_TEST_P(Haar, FaceDetect)
 {
-    MemStorage storage(cvCreateMemStorage(0));
-    CvSeq *_objects;
-    _objects = cascade.oclHaarDetectObjects(d_img, storage, 1.1, 3,
-                                            flags, Size(30, 30), Size(0, 0));
-    vector<CvAvgComp> vecAvgComp;
-    Seq<CvAvgComp>(_objects).copyTo(vecAvgComp);
-    oclfaces.resize(vecAvgComp.size());
-    std::transform(vecAvgComp.begin(), vecAvgComp.end(), oclfaces.begin(), getRect());
+    cascade.detectMultiScale(d_img, oclfaces,  1.1, 3,
+                                flags,
+                                Size(30, 30), Size(0, 0));
 
     cpucascade.detectMultiScale(img, faces,  1.1, 3,
                                 flags,

platforms/android/service/engine/AndroidManifest.xml

@@ -1,8 +1,8 @@
 <?xml version="1.0" encoding="utf-8"?>
 <manifest xmlns:android="http://schemas.android.com/apk/res/android"
     package="org.opencv.engine"
-    android:versionCode="213@ANDROID_PLATFORM_VERSION_CODE@"
-    android:versionName="2.13" >
+    android:versionCode="214@ANDROID_PLATFORM_VERSION_CODE@"
+    android:versionName="2.14" >
 
     <uses-sdk android:minSdkVersion="@ANDROID_NATIVE_API_LEVEL@" />
     <uses-feature android:name="android.hardware.touchscreen" android:required="false"/>

platforms/android/service/readme.txt

@@ -14,20 +14,20 @@ manually using adb tool:
 
 .. code-block:: sh
 
-    adb install OpenCV-2.4.6-android-sdk/apk/OpenCV_2.4.6_Manager_2.9_<platform>.apk
+    adb install OpenCV-2.4.7-android-sdk/apk/OpenCV_2.4.7_Manager_2.14_<platform>.apk
 
 Use the table below to determine proper OpenCV Manager package for your device:
 
-+------------------------------+--------------+---------------------------------------------------+
-| Hardware Platform            | Android ver. | Package name                                      |
-+==============================+==============+===================================================+
-| armeabi-v7a (ARMv7-A + NEON) |    >= 2.3    | OpenCV_2.4.6_Manager_2.9_armv7a-neon.apk          |
-+------------------------------+--------------+---------------------------------------------------+
-| armeabi-v7a (ARMv7-A + NEON) |     = 2.2    | OpenCV_2.4.6_Manager_2.9_armv7a-neon-android8.apk |
-+------------------------------+--------------+---------------------------------------------------+
-| armeabi (ARMv5, ARMv6)       |    >= 2.3    | OpenCV_2.4.6_Manager_2.9_armeabi.apk              |
-+------------------------------+--------------+---------------------------------------------------+
-| Intel x86                    |    >= 2.3    | OpenCV_2.4.6_Manager_2.9_x86.apk                  |
-+------------------------------+--------------+---------------------------------------------------+
-| MIPS                         |    >= 2.3    | OpenCV_2.4.6_Manager_2.9_mips.apk                 |
-+------------------------------+--------------+---------------------------------------------------+
++------------------------------+--------------+----------------------------------------------------+
+| Hardware Platform            | Android ver. | Package name                                       |
++==============================+==============+====================================================+
+| armeabi-v7a (ARMv7-A + NEON) |    >= 2.3    | OpenCV_2.4.7_Manager_2.14_armv7a-neon.apk          |
++------------------------------+--------------+----------------------------------------------------+
+| armeabi-v7a (ARMv7-A + NEON) |     = 2.2    | OpenCV_2.4.7_Manager_2.14_armv7a-neon-android8.apk |
++------------------------------+--------------+----------------------------------------------------+
+| armeabi (ARMv5, ARMv6)       |    >= 2.3    | OpenCV_2.4.7_Manager_2.14_armeabi.apk              |
++------------------------------+--------------+----------------------------------------------------+
+| Intel x86                    |    >= 2.3    | OpenCV_2.4.7_Manager_2.14_x86.apk                  |
++------------------------------+--------------+----------------------------------------------------+
+| MIPS                         |    >= 2.3    | OpenCV_2.4.7_Manager_2.14_mips.apk                 |
++------------------------------+--------------+----------------------------------------------------+

samples/ocl/adaptive_bilateral_filter.cpp

@@ -12,7 +12,9 @@ int main( int argc, const char** argv )
 {
     const char* keys =
         "{ i | input   |          | specify input image }"
-        "{ k | ksize   |     5    | specify kernel size }"
+        "{ k | ksize   |     11   | specify kernel size }"
+        "{ s | sSpace  |     3    | specify sigma space }"
+        "{ c | sColor  |     30   | specify max color }"
         "{ h | help    | false    | print help message }";
 
     CommandLineParser cmd(argc, argv, keys);
@@ -26,27 +28,36 @@ int main( int argc, const char** argv )
 
     string src_path = cmd.get<string>("i");
     int ks = cmd.get<int>("k");
-    const char * winName[] = {"input", "adaptive bilateral CPU", "adaptive bilateral OpenCL", "bilateralFilter OpenCL"};
+    const char * winName[] = {"input", "ABF OpenCL", "BF OpenCL"};
 
-    Mat src = imread(src_path), abFilterCPU;
+    Mat src = imread(src_path);
     if (src.empty())
     {
         cout << "error read image: " << src_path << endl;
         return EXIT_FAILURE;
     }
 
+    double sigmaSpace = cmd.get<int>("s");
+
+    // sigma for checking pixel values. This is used as is in the "normal" bilateral filter,
+    // and it is used as an upper clamp on the adaptive case.
+    double sigmacolor = cmd.get<int>("c");
+
     ocl::oclMat dsrc(src), dABFilter, dBFilter;
-
     Size ksize(ks, ks);
-    adaptiveBilateralFilter(src,abFilterCPU, ksize, 10);
-    ocl::adaptiveBilateralFilter(dsrc, dABFilter, ksize, 10);
-    ocl::bilateralFilter(dsrc, dBFilter, ks, 30, 9);
 
+    // ksize is the total width/height of neighborhood used to calculate local variance.
+    // sigmaSpace is not a priori related to ksize/2.
+    ocl::adaptiveBilateralFilter(dsrc, dABFilter, ksize, sigmaSpace, sigmacolor);
+    ocl::bilateralFilter(dsrc, dBFilter, ks, sigmacolor, sigmaSpace);
     Mat abFilter = dABFilter, bFilter = dBFilter;
+
+    ocl::finish();
+
     imshow(winName[0], src);
-    imshow(winName[1], abFilterCPU);
-    imshow(winName[2], abFilter);
-    imshow(winName[3], bFilter);
+    imshow(winName[1], abFilter);
+    imshow(winName[2], bFilter);
+
     waitKey();
 
     return EXIT_SUCCESS;

samples/ocl/facedetect.cpp

@@ -41,7 +41,7 @@ static double getTime()
 
 
 static void detect( Mat& img, vector<Rect>& faces,
-             ocl::OclCascadeClassifierBuf& cascade,
+             ocl::OclCascadeClassifier& cascade,
              double scale, bool calTime);
 
 
@@ -87,7 +87,7 @@ int main( int argc, const char** argv )
     outputName = cmd.get<string>("o");
     string cascadeName = cmd.get<string>("t");
     double scale = cmd.get<double>("c");
-    ocl::OclCascadeClassifierBuf cascade;
+    ocl::OclCascadeClassifier cascade;
     CascadeClassifier  cpu_cascade;
 
     if( !cascade.load( cascadeName ) || !cpu_cascade.load(cascadeName) )
@@ -180,7 +180,7 @@ int main( int argc, const char** argv )
 }
 
 void detect( Mat& img, vector<Rect>& faces,
-             ocl::OclCascadeClassifierBuf& cascade,
+             ocl::OclCascadeClassifier& cascade,
              double scale, bool calTime)
 {
     ocl::oclMat image(img);