Merge pull request #1103 from apavlenko:ver_2.4.6.1

- set init value for `numsdev` to prevent use of uninitialized value
- stop use of 'workdir' and files from samples
- forcing use of 'opencv_extra' instead
Note: set OPENCV_TEST_DATA_PATH to full path to 'opencv_extra/testdata' (gitolite@code.opencv.org:opencv_extra.git) before running the test!

3rdparty/tbb/CMakeLists.txt

@@ -5,14 +5,23 @@ if (WIN32 AND NOT ARM)
   message(FATAL_ERROR "BUILD_TBB option supports Windows on ARM only!\nUse regular official TBB build instead of the BUILD_TBB option!")
 endif()
 
-# 4.1 update 4 - works fine
-set(tbb_ver "tbb41_20130613oss")
-set(tbb_url "http://threadingbuildingblocks.org/sites/default/files/software_releases/source/tbb41_20130613oss_src.tgz")
-set(tbb_md5 "108c8c1e481b0aaea61878289eb28b6a")
-set(tbb_version_file "version_string.ver")
-ocv_warnings_disable(CMAKE_CXX_FLAGS -Wshadow -Wunused-parameter)
+if (WIN32 AND ARM)
+  # 4.1 update 4 - The first release that supports Windows RT. Hangs on some Android devices
+  set(tbb_ver "tbb41_20130613oss")
+  set(tbb_url "http://threadingbuildingblocks.org/sites/default/files/software_releases/source/tbb41_20130613oss_src.tgz")
+  set(tbb_md5 "108c8c1e481b0aaea61878289eb28b6a")
+  set(tbb_version_file "version_string.ver")
+  ocv_warnings_disable(CMAKE_CXX_FLAGS -Wshadow -Wunused-parameter)
+else()
+  # 4.1 update 2 - works fine
+  set(tbb_ver "tbb41_20130116oss")
+  set(tbb_url "http://threadingbuildingblocks.org/sites/default/files/software_releases/source/tbb41_20130116oss_src.tgz")
+  set(tbb_md5 "3809790e1001a1b32d59c9fee590ee85")
+  set(tbb_version_file "version_string.ver")
+  ocv_warnings_disable(CMAKE_CXX_FLAGS -Wshadow)
+endif()
 
-# 4.1 update 3 dev - works fine
+# 4.1 update 3 dev - Hangs on some Android devices
 #set(tbb_ver "tbb41_20130401oss")
 #set(tbb_url "http://threadingbuildingblocks.org/sites/default/files/software_releases/source/tbb41_20130401oss_src.tgz")
 #set(tbb_md5 "f2f591a0d2ca8f801e221ce7d9ea84bb")

doc/tutorials/introduction/android_binary_package/O4A_SDK.rst

@@ -51,7 +51,7 @@ The structure of package contents looks as follows:
     OpenCV-2.4.6-android-sdk
     |_ apk
     |   |_ OpenCV_2.4.6_binary_pack_armv7a.apk
-    |   |_ OpenCV_2.4.6_Manager_2.8_XXX.apk
+    |   |_ OpenCV_2.4.6_Manager_2.9_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.6_Manager_2.8_armv7a-neon.apk
+    <Android SDK path>/platform-tools/adb install <OpenCV4Android SDK path>/apk/OpenCV_2.4.6_Manager_2.9_armv7a-neon.apk
 
   .. note:: ``armeabi``, ``armv7a-neon``, ``arm7a-neon-android8``, ``mips`` and ``x86`` stand for
             platform targets:

modules/core/include/opencv2/core/version.hpp

@@ -50,7 +50,7 @@
 #define CV_VERSION_EPOCH    2
 #define CV_VERSION_MAJOR    4
 #define CV_VERSION_MINOR    6
-#define CV_VERSION_REVISION 0
+#define CV_VERSION_REVISION 1
 
 #define CVAUX_STR_EXP(__A)  #__A
 #define CVAUX_STR(__A)      CVAUX_STR_EXP(__A)

modules/features2d/include/opencv2/features2d/features2d.hpp

@@ -589,11 +589,11 @@ protected:
 };
 
 
-class CV_EXPORTS GFTTDetector : public FeatureDetector
+class CV_EXPORTS_W GFTTDetector : public FeatureDetector
 {
 public:
-    GFTTDetector( int maxCorners=1000, double qualityLevel=0.01, double minDistance=1,
-                  int blockSize=3, bool useHarrisDetector=false, double k=0.04 );
+    CV_WRAP GFTTDetector( int maxCorners=1000, double qualityLevel=0.01, double minDistance=1,
+                          int blockSize=3, bool useHarrisDetector=false, double k=0.04 );
     AlgorithmInfo* info() const;
 
 protected:

modules/highgui/src/cap.cpp

@@ -220,8 +220,8 @@ CV_IMPL CvCapture * cvCreateCameraCapture (int index)
                 return capture;
             break;
 #endif
-#ifdef HAVE_VFW
         case CV_CAP_VFW:
+#ifdef HAVE_VFW
             capture = cvCreateCameraCapture_VFW (index);
             if (capture)
                 return capture;

modules/highgui/src/cap_qtkit.mm

@@ -277,11 +277,8 @@ bool CvCaptureCAM::grabFrame(double timeOut) {
     double sleepTime = 0.005;
     double total = 0;
 
-    NSDate *loopUntil = [NSDate dateWithTimeIntervalSinceNow:sleepTime];
-    while (![capture updateImage] && (total += sleepTime)<=timeOut &&
-           [[NSRunLoop currentRunLoop] runMode: NSDefaultRunLoopMode
-                                    beforeDate:loopUntil])
-        loopUntil = [NSDate dateWithTimeIntervalSinceNow:sleepTime];
+    while (![capture updateImage] && (total += sleepTime)<=timeOut)
+        usleep((int)(sleepTime*1000));
 
     [localpool drain];
 

modules/java/generator/src/java/android+JavaCameraView.java

@@ -38,7 +38,7 @@ public class JavaCameraView extends CameraBridgeViewBase implements PreviewCallb
     private boolean mStopThread;
 
     protected Camera mCamera;
-    protected JavaCameraFrame mCameraFrame;
+    protected JavaCameraFrame[] mCameraFrame;
     private SurfaceTexture mSurfaceTexture;
 
     public static class JavaCameraSizeAccessor implements ListItemAccessor {
@@ -180,7 +180,9 @@ public class JavaCameraView extends CameraBridgeViewBase implements PreviewCallb
 
                     AllocateCache();
 
-                    mCameraFrame = new JavaCameraFrame(mFrameChain[mChainIdx], mFrameWidth, mFrameHeight);
+                    mCameraFrame = new JavaCameraFrame[2];
+                    mCameraFrame[0] = new JavaCameraFrame(mFrameChain[0], mFrameWidth, mFrameHeight);
+                    mCameraFrame[1] = new JavaCameraFrame(mFrameChain[1], mFrameWidth, mFrameHeight);
 
                     if (Build.VERSION.SDK_INT >= Build.VERSION_CODES.HONEYCOMB) {
                         mSurfaceTexture = new SurfaceTexture(MAGIC_TEXTURE_ID);
@@ -216,8 +218,10 @@ public class JavaCameraView extends CameraBridgeViewBase implements PreviewCallb
                 mFrameChain[0].release();
                 mFrameChain[1].release();
             }
-            if (mCameraFrame != null)
-                mCameraFrame.release();
+            if (mCameraFrame != null) {
+                mCameraFrame[0].release();
+                mCameraFrame[1].release();
+            }
         }
     }
 
@@ -318,7 +322,7 @@ public class JavaCameraView extends CameraBridgeViewBase implements PreviewCallb
 
                 if (!mStopThread) {
                     if (!mFrameChain[mChainIdx].empty())
-                        deliverAndDrawFrame(mCameraFrame);
+                        deliverAndDrawFrame(mCameraFrame[mChainIdx]);
                     mChainIdx = 1 - mChainIdx;
                 }
             } while (!mStopThread);

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

@@ -211,8 +211,8 @@ namespace cv
         {
             flags &= roi.width < m.cols ? ~Mat::CONTINUOUS_FLAG : -1;
             offset += roi.y * step + roi.x * elemSize();
-            CV_Assert( 0 <= roi.x && 0 <= roi.width && roi.x + roi.width <= m.cols &&
-                       0 <= roi.y && 0 <= roi.height && roi.y + roi.height <= m.rows );
+            CV_Assert( 0 <= roi.x && 0 <= roi.width && roi.x + roi.width <= m.wholecols &&
+                       0 <= roi.y && 0 <= roi.height && roi.y + roi.height <= m.wholerows );
             if( refcount )
                 CV_XADD(refcount, 1);
             if( rows <= 0 || cols <= 0 )

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

@@ -1395,6 +1395,45 @@ namespace cv
             oclMat vPyr_[2];
             bool isDeviceArch11_;
         };
+
+        class CV_EXPORTS FarnebackOpticalFlow
+        {
+        public:
+            FarnebackOpticalFlow();
+
+            int numLevels;
+            double pyrScale;
+            bool fastPyramids;
+            int winSize;
+            int numIters;
+            int polyN;
+            double polySigma;
+            int flags;
+
+            void operator ()(const oclMat &frame0, const oclMat &frame1, oclMat &flowx, oclMat &flowy);
+
+            void releaseMemory();
+
+        private:
+            void prepareGaussian(
+                int n, double sigma, float *g, float *xg, float *xxg,
+                double &ig11, double &ig03, double &ig33, double &ig55);
+
+            void setPolynomialExpansionConsts(int n, double sigma);
+
+            void updateFlow_boxFilter(
+                const oclMat& R0, const oclMat& R1, oclMat& flowx, oclMat &flowy,
+                oclMat& M, oclMat &bufM, int blockSize, bool updateMatrices);
+
+            void updateFlow_gaussianBlur(
+                const oclMat& R0, const oclMat& R1, oclMat& flowx, oclMat& flowy,
+                oclMat& M, oclMat &bufM, int blockSize, bool updateMatrices);
+
+            oclMat frames_[2];
+            oclMat pyrLevel_[2], M_, bufM_, R_[2], blurredFrame_[2];
+            std::vector<oclMat> pyramid0_, pyramid1_;
+        };
+
         //////////////// build warping maps ////////////////////
         //! builds plane warping maps
         CV_EXPORTS void buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat &R, const Mat &T, float scale, oclMat &map_x, oclMat &map_y);

modules/ocl/perf/perf_opticalflow.cpp

@@ -136,11 +136,13 @@ PERFTEST(PyrLKOpticalFlow)
             size_t mismatch = 0;
             for (int i = 0; i < (int)nextPts.size(); ++i)
             {
-                if(status[i] != ocl_status.at<unsigned char>(0, i)){
+                if(status[i] != ocl_status.at<unsigned char>(0, i))
+                {
                     mismatch++;
                     continue;
                 }
-                if(status[i]){
+                if(status[i])
+                {
                     Point2f gpu_rst = ocl_nextPts.at<Point2f>(0, i);
                     Point2f cpu_rst = nextPts[i];
                     if(fabs(gpu_rst.x - cpu_rst.x) >= 1. || fabs(gpu_rst.y - cpu_rst.y) >= 1.)
@@ -193,24 +195,24 @@ PERFTEST(tvl1flow)
     WARMUP_ON;
     d_alg(d0, d1, d_flowx, d_flowy);
     WARMUP_OFF;
-/*
-    double diff1 = 0.0, diff2 = 0.0;
-    if(ExceptedMatSimilar(gold[0], cv::Mat(d_flowx), 3e-3, diff1) == 1
-        &&ExceptedMatSimilar(gold[1], cv::Mat(d_flowy), 3e-3, diff2) == 1)
-        TestSystem::instance().setAccurate(1);
-    else
-        TestSystem::instance().setAccurate(0);
+    /*
+        double diff1 = 0.0, diff2 = 0.0;
+        if(ExceptedMatSimilar(gold[0], cv::Mat(d_flowx), 3e-3, diff1) == 1
+            &&ExceptedMatSimilar(gold[1], cv::Mat(d_flowy), 3e-3, diff2) == 1)
+            TestSystem::instance().setAccurate(1);
+        else
+            TestSystem::instance().setAccurate(0);
 
-    TestSystem::instance().setDiff(diff1);
-    TestSystem::instance().setDiff(diff2);
-*/
+        TestSystem::instance().setDiff(diff1);
+        TestSystem::instance().setDiff(diff2);
+    */
 
 
     GPU_ON;
     d_alg(d0, d1, d_flowx, d_flowy);
     d_alg.collectGarbage();
     GPU_OFF;
-    
+
 
     cv::Mat flowx, flowy;
 
@@ -225,4 +227,130 @@ PERFTEST(tvl1flow)
 
     TestSystem::instance().ExceptedMatSimilar(gold[0], flowx, 3e-3);
     TestSystem::instance().ExceptedMatSimilar(gold[1], flowy, 3e-3);
-}
+}
+
+///////////// FarnebackOpticalFlow ////////////////////////
+PERFTEST(FarnebackOpticalFlow)
+{
+    cv::Mat frame0 = imread("rubberwhale1.png", cv::IMREAD_GRAYSCALE);
+    ASSERT_FALSE(frame0.empty());
+
+    cv::Mat frame1 = imread("rubberwhale2.png", cv::IMREAD_GRAYSCALE);
+    ASSERT_FALSE(frame1.empty());
+
+    cv::ocl::oclMat d_frame0(frame0), d_frame1(frame1);
+
+    int polyNs[2] = { 5, 7 };
+    double polySigmas[2] = { 1.1, 1.5 };
+    int farneFlags[2] = { 0, cv::OPTFLOW_FARNEBACK_GAUSSIAN };
+    bool UseInitFlows[2] = { false, true };
+    double pyrScale = 0.5;
+
+    string farneFlagStrs[2] = { "BoxFilter", "GaussianBlur" };
+    string useInitFlowStrs[2] = { "", "UseInitFlow" };
+
+    for ( int i = 0; i < 2; ++i)
+    {
+        int polyN = polyNs[i];
+        double polySigma = polySigmas[i];
+
+        for ( int j = 0; j < 2; ++j)
+        {
+            int flags = farneFlags[j];
+
+            for ( int k = 0; k < 2; ++k)
+            {
+                bool useInitFlow = UseInitFlows[k];
+                SUBTEST << "polyN(" << polyN << "); " << farneFlagStrs[j] << "; " << useInitFlowStrs[k];
+
+                cv::ocl::FarnebackOpticalFlow farn;
+                farn.pyrScale = pyrScale;
+                farn.polyN = polyN;
+                farn.polySigma = polySigma;
+                farn.flags = flags;
+
+                cv::ocl::oclMat d_flowx, d_flowy;
+                cv::Mat flow, flowBuf, flowxBuf, flowyBuf;
+
+                WARMUP_ON;
+                farn(d_frame0, d_frame1, d_flowx, d_flowy);
+
+                if (useInitFlow)
+                {
+                    cv::Mat flowxy[] = {cv::Mat(d_flowx), cv::Mat(d_flowy)};
+                    cv::merge(flowxy, 2, flow);
+                    flow.copyTo(flowBuf);
+                    flowxy[0].copyTo(flowxBuf);
+                    flowxy[1].copyTo(flowyBuf);
+
+                    farn.flags |= cv::OPTFLOW_USE_INITIAL_FLOW;
+                    farn(d_frame0, d_frame1, d_flowx, d_flowy);
+                }
+                WARMUP_OFF;
+
+                cv::calcOpticalFlowFarneback(
+                    frame0, frame1, flow, farn.pyrScale, farn.numLevels, farn.winSize,
+                    farn.numIters, farn.polyN, farn.polySigma, farn.flags);
+
+                std::vector<cv::Mat> flowxy;
+                cv::split(flow, flowxy);
+
+                Mat md_flowx = cv::Mat(d_flowx);
+                Mat md_flowy = cv::Mat(d_flowy);
+                TestSystem::instance().ExceptedMatSimilar(flowxy[0], md_flowx, 0.1);
+                TestSystem::instance().ExceptedMatSimilar(flowxy[1], md_flowy, 0.1);
+
+                if (useInitFlow)
+                {
+                    cv::Mat flowx, flowy;
+                    farn.flags = (flags | cv::OPTFLOW_USE_INITIAL_FLOW);
+
+                    CPU_ON;
+                    cv::calcOpticalFlowFarneback(
+                        frame0, frame1, flowBuf, farn.pyrScale, farn.numLevels, farn.winSize,
+                        farn.numIters, farn.polyN, farn.polySigma, farn.flags);
+                    CPU_OFF;
+
+                    GPU_ON;
+                    farn(d_frame0, d_frame1, d_flowx, d_flowy);
+                    GPU_OFF;
+
+                    GPU_FULL_ON;
+                    d_frame0.upload(frame0);
+                    d_frame1.upload(frame1);
+                    d_flowx.upload(flowxBuf);
+                    d_flowy.upload(flowyBuf);
+                    farn(d_frame0, d_frame1, d_flowx, d_flowy);
+                    d_flowx.download(flowx);
+                    d_flowy.download(flowy);
+                    GPU_FULL_OFF;
+                }
+                else
+                {
+                    cv::Mat flow, flowx, flowy;
+                    cv::ocl::oclMat d_flowx, d_flowy;
+
+                    farn.flags = flags;
+
+                    CPU_ON;
+                    cv::calcOpticalFlowFarneback(
+                        frame0, frame1, flow, farn.pyrScale, farn.numLevels, farn.winSize,
+                        farn.numIters, farn.polyN, farn.polySigma, farn.flags);
+                    CPU_OFF;
+
+                    GPU_ON;
+                    farn(d_frame0, d_frame1, d_flowx, d_flowy);
+                    GPU_OFF;
+
+                    GPU_FULL_ON;
+                    d_frame0.upload(frame0);
+                    d_frame1.upload(frame1);
+                    farn(d_frame0, d_frame1, d_flowx, d_flowy);
+                    d_flowx.download(flowx);
+                    d_flowy.download(flowy);
+                    GPU_FULL_OFF;
+                }
+            }
+        }
+    }
+}

modules/ocl/src/hog.cpp

@@ -1816,8 +1816,14 @@ void cv::ocl::device::hog::normalize_hists(int nbins,
         openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads, 
                              localThreads, args, -1, -1, "-D CPU");
     else
+    {
+        cl_kernel kernel = openCLGetKernelFromSource(clCxt, &objdetect_hog, kernelName);
+        int wave_size = queryDeviceInfo<WAVEFRONT_SIZE, int>(kernel);
+        char opt[32] = {0};
+        sprintf(opt, "-D WAVE_SIZE=%d", wave_size);
         openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads, 
-                             localThreads, args, -1, -1);
+                             localThreads, args, -1, -1, opt);
+    }
 }
 
 void cv::ocl::device::hog::classify_hists(int win_height, int win_width, 
@@ -1879,8 +1885,14 @@ void cv::ocl::device::hog::classify_hists(int win_height, int win_width,
         openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads, 
                              localThreads, args, -1, -1, "-D CPU");
     else
+    {
+        cl_kernel kernel = openCLGetKernelFromSource(clCxt, &objdetect_hog, kernelName);
+        int wave_size = queryDeviceInfo<WAVEFRONT_SIZE, int>(kernel);
+        char opt[32] = {0};
+        sprintf(opt, "-D WAVE_SIZE=%d", wave_size);
         openCLExecuteKernel(clCxt, &objdetect_hog, kernelName, globalThreads, 
-                             localThreads, args, -1, -1);
+                             localThreads, args, -1, -1, opt);
+    }
 }
 
 void cv::ocl::device::hog::extract_descrs_by_rows(int win_height, int win_width, 

modules/ocl/src/initialization.cpp

@@ -123,9 +123,6 @@ namespace cv
             codeCache.clear();
             cacheSize = 0;
         }
-
-        // not to be exported to dynamic lib
-        void setBinaryDiskCacheImpl(int mode, String path, Info::Impl * impl);
         struct Info::Impl
         {
             cl_platform_id oclplatform;
@@ -143,9 +140,6 @@ namespace cv
             char extra_options[512];
             int  double_support;
             int unified_memory; //1 means integrated GPU, otherwise this value is 0
-            bool enable_disk_cache; 
-            bool update_disk_cache;
-            string binpath;
             int refcounter;
 
             Impl();
@@ -173,6 +167,16 @@ namespace cv
             void releaseResources();
         };
 
+        // global variables to hold binary cache properties
+        static int enable_disk_cache = 
+#ifdef _DEBUG
+            false;
+#else
+            true;
+#endif
+        static int update_disk_cache = false;
+        static String binpath = "";
+
         Info::Impl::Impl()
             :oclplatform(0),
             oclcontext(0),
@@ -183,13 +187,9 @@ namespace cv
             maxComputeUnits(0),
             double_support(0),
             unified_memory(0),
-            enable_disk_cache(false),
-            update_disk_cache(false),
-            binpath("./"),
             refcounter(1)
         {
             memset(extra_options, 0, 512);
-            setBinaryDiskCacheImpl(CACHE_RELEASE, String("./"), this);
         }
 
         void Info::Impl::releaseResources()
@@ -198,7 +198,8 @@ namespace cv
 
             if(clCmdQueue)
             {
-                openCLSafeCall(clReleaseCommandQueue(clCmdQueue));
+                //temporarily disable command queue release as it causes program hang at exit
+                //openCLSafeCall(clReleaseCommandQueue(clCmdQueue));
                 clCmdQueue = 0;
             }
 
@@ -318,7 +319,7 @@ namespace cv
             char clVersion[256];
             for (unsigned i = 0; i < numPlatforms; ++i)
             {
-                cl_uint numsdev;
+                cl_uint numsdev = 0;
                 cl_int status = clGetDeviceIDs(platforms[i], devicetype, 0, NULL, &numsdev);
                 if(status != CL_DEVICE_NOT_FOUND)
                     openCLVerifyCall(status);
@@ -504,29 +505,30 @@ namespace cv
             return openCLGetKernelFromSource(clCxt, source, kernelName, NULL);
         }
 
-        void setBinaryDiskCacheImpl(int mode, String path, Info::Impl * impl)
+        void setBinaryDiskCache(int mode, String path)
         {
-            impl->update_disk_cache = (mode & CACHE_UPDATE) == CACHE_UPDATE;
-            impl->enable_disk_cache = 
+            if(mode == CACHE_NONE)
+            {
+                update_disk_cache = 0;
+                enable_disk_cache = 0;
+                return;
+            }
+            update_disk_cache |= (mode & CACHE_UPDATE) == CACHE_UPDATE;
+            enable_disk_cache |= 
 #ifdef _DEBUG 
                 (mode & CACHE_DEBUG)   == CACHE_DEBUG;
 #else
                 (mode & CACHE_RELEASE) == CACHE_RELEASE;
 #endif
-            if(impl->enable_disk_cache && !path.empty())
+            if(enable_disk_cache && !path.empty())
             {
-                impl->binpath = path;
+                binpath = path;
             }
         }
-        void setBinaryDiskCache(int mode, cv::String path)
-        {
-            setBinaryDiskCacheImpl(mode, path, Context::getContext()->impl);
-        }
 
         void setBinpath(const char *path)
         {
-            Context *clcxt = Context::getContext();
-            clcxt->impl->binpath = path;
+            binpath = path;
         }
 
         int savetofile(const Context*,  cl_program &program, const char *fileName)
@@ -594,15 +596,15 @@ namespace cv
                     strcat(all_build_options, build_options);
                 if(all_build_options != NULL)
                 {
-                    filename = clCxt->impl->binpath  + kernelName + "_" + clCxt->impl->devName[clCxt->impl->devnum] + all_build_options + ".clb";
+                    filename = binpath + kernelName + "_" + clCxt->impl->devName[clCxt->impl->devnum] + all_build_options + ".clb";
                 }
                 else
                 {
-                    filename = clCxt->impl->binpath  + kernelName + "_" + clCxt->impl->devName[clCxt->impl->devnum] + ".clb";
+                    filename = binpath + kernelName + "_" + clCxt->impl->devName[clCxt->impl->devnum] + ".clb";
                 }
 
-                FILE *fp = clCxt->impl->enable_disk_cache ? fopen(filename.c_str(), "rb") : NULL;
-                if(fp == NULL || clCxt->impl->update_disk_cache)
+                FILE *fp = enable_disk_cache ? fopen(filename.c_str(), "rb") : NULL;
+                if(fp == NULL || update_disk_cache)
                 {
                     if(fp != NULL)
                         fclose(fp);
@@ -611,7 +613,7 @@ namespace cv
                                   clCxt->impl->oclcontext, 1, source, NULL, &status);
                     openCLVerifyCall(status);
                     status = clBuildProgram(program, 1, &(clCxt->impl->devices[clCxt->impl->devnum]), all_build_options, NULL, NULL);
-                    if(status == CL_SUCCESS && clCxt->impl->enable_disk_cache)
+                    if(status == CL_SUCCESS && enable_disk_cache)
                         savetofile(clCxt, program, filename.c_str());
                 }
                 else
@@ -1075,26 +1077,3 @@ namespace cv
     }//namespace ocl
 
 }//namespace cv
-
-#if defined BUILD_SHARED_LIBS && defined CVAPI_EXPORTS && defined WIN32 && !defined WINCE
-#include <windows.h>
-BOOL WINAPI DllMain( HINSTANCE, DWORD  fdwReason, LPVOID );
-
-BOOL WINAPI DllMain( HINSTANCE, DWORD  fdwReason, LPVOID )
-{
-    if( fdwReason == DLL_PROCESS_DETACH )
-    {
-        // application hangs if call clReleaseCommandQueue here, so release context only
-        // without context release application hangs as well
-        context_tear_down = 1;
-        Context* cv_ctx = Context::getContext();
-        if(cv_ctx)
-        {
-            cl_context ctx = cv_ctx->impl->oclcontext;
-            if(ctx)
-                openCLSafeCall(clReleaseContext(ctx));
-        }
-    }
-    return TRUE;
-}
-#endif

modules/ocl/src/opencl/objdetect_hog.cl

@@ -133,7 +133,9 @@ __kernel void compute_hists_lut_kernel(
             final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] = 
                 hist_[0] + hist_[1] + hist_[2];
     }
+#ifdef CPU
     barrier(CLK_LOCAL_MEM_FENCE);
+#endif
 
     int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 12 + cell_thread_x;
     if ((tid < cblock_hist_size) && (gid < blocks_total))
@@ -225,8 +227,9 @@ __kernel void compute_hists_kernel(
             final_hist[(cell_x * 2 + cell_y) * cnbins + bin_id] = 
                 hist_[0] + hist_[1] + hist_[2];
     }
+#ifdef CPU
     barrier(CLK_LOCAL_MEM_FENCE);
-
+#endif
     int tid = (cell_y * CELLS_PER_BLOCK_Y + cell_x) * 12 + cell_thread_x;
     if ((tid < cblock_hist_size) && (gid < blocks_total))
     {
@@ -318,6 +321,10 @@ float reduce_smem(volatile __local float* smem, int size)
     if (tid < 32)
     {
         if (size >= 64) smem[tid] = sum = sum + smem[tid + 32];
+#if WAVE_SIZE < 32
+    } barrier(CLK_LOCAL_MEM_FENCE);
+    if (tid < 16) {
+#endif
         if (size >= 32) smem[tid] = sum = sum + smem[tid + 16];
         if (size >= 16) smem[tid] = sum = sum + smem[tid + 8];
         if (size >= 8) smem[tid] = sum = sum + smem[tid + 4];
@@ -418,6 +425,9 @@ __kernel void classify_hists_180_kernel(
     {
         smem[tid] = product = product + smem[tid + 32];
     }
+#if WAVE_SIZE < 32
+    barrier(CLK_LOCAL_MEM_FENCE);
+#endif
     if (tid < 16)
     {
         smem[tid] = product = product + smem[tid + 16];
@@ -487,6 +497,10 @@ __kernel void classify_hists_252_kernel(
     if (tid < 32)
     {      
         smem[tid] = product = product + smem[tid + 32];
+#if WAVE_SIZE < 32
+    } barrier(CLK_LOCAL_MEM_FENCE);
+    if (tid < 16) {
+#endif
         smem[tid] = product = product + smem[tid + 16];
         smem[tid] = product = product + smem[tid + 8];
         smem[tid] = product = product + smem[tid + 4];
@@ -553,6 +567,10 @@ __kernel void classify_hists_kernel(
     if (tid < 32)
     {       
         smem[tid] = product = product + smem[tid + 32];
+#if WAVE_SIZE < 32
+    } barrier(CLK_LOCAL_MEM_FENCE);
+    if (tid < 16) {
+#endif
         smem[tid] = product = product + smem[tid + 16];
         smem[tid] = product = product + smem[tid + 8];
         smem[tid] = product = product + smem[tid + 4];

modules/ocl/src/opencl/optical_flow_farneback.cl

@@ -0,0 +1,450 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//    Sen Liu, swjtuls1987@126.com
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other oclMaterials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors as is and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+
+#define tx  get_local_id(0)
+#define ty  get_local_id(1)
+#define bx  get_group_id(0)
+#define bdx get_local_size(0)
+
+#define BORDER_SIZE 5
+#define MAX_KSIZE_HALF 100
+
+#ifndef polyN
+#define polyN 5
+#endif
+
+__kernel void polynomialExpansion(__global float * dst,
+                                  __global __const float * src,
+                                  __global __const float * c_g,
+                                  __global __const float * c_xg,
+                                  __global __const float * c_xxg,
+                                  __local float * smem,
+                                  const float4 ig,
+                                  const int height, const int width,
+                                  int dstStep, int srcStep)
+{
+    const int y = get_global_id(1);
+    const int x = bx * (bdx - 2*polyN) + tx - polyN;
+
+    dstStep /= sizeof(*dst);
+    srcStep /= sizeof(*src);
+
+    int xWarped;
+    __local float *row = smem + tx;
+
+    if (y < height && y >= 0)
+    {
+        xWarped = min(max(x, 0), width - 1);
+
+        row[0] = src[mad24(y, srcStep, xWarped)] * c_g[0];
+        row[bdx] = 0.f;
+        row[2*bdx] = 0.f;
+
+#pragma unroll
+        for (int k = 1; k <= polyN; ++k)
+        {
+            float t0 = src[mad24(max(y - k, 0), srcStep, xWarped)];
+            float t1 = src[mad24(min(y + k, height - 1), srcStep, xWarped)];
+
+            row[0] += c_g[k] * (t0 + t1);
+            row[bdx] += c_xg[k] * (t1 - t0);
+            row[2*bdx] += c_xxg[k] * (t0 + t1);
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (y < height && y >= 0 && tx >= polyN && tx + polyN < bdx && x < width)
+    {
+        float b1 = c_g[0] * row[0];
+        float b3 = c_g[0] * row[bdx];
+        float b5 = c_g[0] * row[2*bdx];
+        float b2 = 0, b4 = 0, b6 = 0;
+
+#pragma unroll
+        for (int k = 1; k <= polyN; ++k)
+        {
+            b1 += (row[k] + row[-k]) * c_g[k];
+            b4 += (row[k] + row[-k]) * c_xxg[k];
+            b2 += (row[k] - row[-k]) * c_xg[k];
+            b3 += (row[k + bdx] + row[-k + bdx]) * c_g[k];
+            b6 += (row[k + bdx] - row[-k + bdx]) * c_xg[k];
+            b5 += (row[k + 2*bdx] + row[-k + 2*bdx]) * c_g[k];
+        }
+
+        dst[mad24(y, dstStep, xWarped)] = b3*ig.s0;
+        dst[mad24(height + y, dstStep, xWarped)] = b2*ig.s0;
+        dst[mad24(2*height + y, dstStep, xWarped)] = b1*ig.s1 + b5*ig.s2;
+        dst[mad24(3*height + y, dstStep, xWarped)] = b1*ig.s1 + b4*ig.s2;
+        dst[mad24(4*height + y, dstStep, xWarped)] = b6*ig.s3;
+    }
+}
+
+inline int idx_row_low(const int y, const int last_row)
+{
+    return abs(y) % (last_row + 1);
+}
+
+inline int idx_row_high(const int y, const int last_row)
+{
+    return abs(last_row - abs(last_row - y)) % (last_row + 1);
+}
+
+inline int idx_row(const int y, const int last_row)
+{
+    return idx_row_low(idx_row_high(y, last_row), last_row);
+}
+
+inline int idx_col_low(const int x, const int last_col)
+{
+    return abs(x) % (last_col + 1);
+}
+
+inline int idx_col_high(const int x, const int last_col)
+{
+    return abs(last_col - abs(last_col - x)) % (last_col + 1);
+}
+
+inline int idx_col(const int x, const int last_col)
+{
+    return idx_col_low(idx_col_high(x, last_col), last_col);
+}
+
+__kernel void gaussianBlur(__global float * dst,
+                           __global const float * src,
+                           __global const float * c_gKer,
+                           __local float * smem,
+                           const int height,  const int width,
+                           int dstStep, int srcStep,
+                           const int ksizeHalf)
+{
+    const int y = get_global_id(1);
+    const int x = get_global_id(0);
+
+    dstStep /= sizeof(*dst);
+    srcStep /= sizeof(*src);
+
+    __local float *row = smem + ty * (bdx + 2*ksizeHalf);
+
+    if (y < height)
+    {
+        // Vertical pass
+        for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx)
+        {
+            int xExt = (int)(bx * bdx) + i - ksizeHalf;
+            xExt = idx_col(xExt, width - 1);
+            row[i] = src[mad24(y, srcStep, xExt)] * c_gKer[0];
+            for (int j = 1; j <= ksizeHalf; ++j)
+                row[i] += (src[mad24(idx_row_low(y - j, height - 1), srcStep, xExt)]
+                           + src[mad24(idx_row_high(y + j, height - 1), srcStep, xExt)]) * c_gKer[j];
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (y < height && y >= 0 && x < width && x >= 0)
+    {
+        // Horizontal pass
+        row += tx + ksizeHalf;
+        float res = row[0] * c_gKer[0];
+        for (int i = 1; i <= ksizeHalf; ++i)
+            res += (row[-i] + row[i]) * c_gKer[i];
+
+        dst[mad24(y, dstStep, x)] = res;
+    }
+}
+
+__constant float c_border[BORDER_SIZE + 1] = { 0.14f, 0.14f, 0.4472f, 0.4472f, 0.4472f, 1.f };
+
+__kernel void updateMatrices(__global float * M,
+                             __global const float * flowx, __global const float * flowy,
+                             __global const float * R0, __global const float * R1,
+                             const int height, const int width,
+                             int mStep, int xStep,  int yStep, int R0Step, int R1Step)
+{
+    const int y = get_global_id(1);
+    const int x = get_global_id(0);
+
+    mStep /= sizeof(*M);
+    xStep /= sizeof(*flowx);
+    yStep /= sizeof(*flowy);
+    R0Step /= sizeof(*R0);
+    R1Step /= sizeof(*R1);
+
+    if (y < height && y >= 0 && x < width && x >= 0)
+    {
+        float dx = flowx[mad24(y, xStep, x)];
+        float dy = flowy[mad24(y, yStep, x)];
+        float fx = x + dx;
+        float fy = y + dy;
+
+        int x1 = convert_int(floor(fx));
+        int y1 = convert_int(floor(fy));
+        fx -= x1;
+        fy -= y1;
+
+        float r2, r3, r4, r5, r6;
+
+        if (x1 >= 0 && y1 >= 0 && x1 < width - 1 && y1 < height - 1)
+        {
+            float a00 = (1.f - fx) * (1.f - fy);
+            float a01 = fx * (1.f - fy);
+            float a10 = (1.f - fx) * fy;
+            float a11 = fx * fy;
+
+            r2 = a00 * R1[mad24(y1, R1Step, x1)] +
+                 a01 * R1[mad24(y1, R1Step, x1 + 1)] +
+                 a10 * R1[mad24(y1 + 1, R1Step, x1)] +
+                 a11 * R1[mad24(y1 + 1, R1Step, x1 + 1)];
+
+            r3 = a00 * R1[mad24(height + y1, R1Step, x1)] +
+                 a01 * R1[mad24(height + y1, R1Step, x1 + 1)] +
+                 a10 * R1[mad24(height + y1 + 1, R1Step, x1)] +
+                 a11 * R1[mad24(height + y1 + 1, R1Step, x1 + 1)];
+
+            r4 = a00 * R1[mad24(2*height + y1, R1Step, x1)] +
+                 a01 * R1[mad24(2*height + y1, R1Step, x1 + 1)] +
+                 a10 * R1[mad24(2*height + y1 + 1, R1Step, x1)] +
+                 a11 * R1[mad24(2*height + y1 + 1, R1Step, x1 + 1)];
+
+            r5 = a00 * R1[mad24(3*height + y1, R1Step, x1)] +
+                 a01 * R1[mad24(3*height + y1, R1Step, x1 + 1)] +
+                 a10 * R1[mad24(3*height + y1 + 1, R1Step, x1)] +
+                 a11 * R1[mad24(3*height + y1 + 1, R1Step, x1 + 1)];
+
+            r6 = a00 * R1[mad24(4*height + y1, R1Step, x1)] +
+                 a01 * R1[mad24(4*height + y1, R1Step, x1 + 1)] +
+                 a10 * R1[mad24(4*height + y1 + 1, R1Step, x1)] +
+                 a11 * R1[mad24(4*height + y1 + 1, R1Step, x1 + 1)];
+
+            r4 = (R0[mad24(2*height + y, R0Step, x)] + r4) * 0.5f;
+            r5 = (R0[mad24(3*height + y, R0Step, x)] + r5) * 0.5f;
+            r6 = (R0[mad24(4*height + y, R0Step, x)] + r6) * 0.25f;
+        }
+        else
+        {
+            r2 = r3 = 0.f;
+            r4 = R0[mad24(2*height + y, R0Step, x)];
+            r5 = R0[mad24(3*height + y, R0Step, x)];
+            r6 = R0[mad24(4*height + y, R0Step, x)] * 0.5f;
+        }
+
+        r2 = (R0[mad24(y, R0Step, x)] - r2) * 0.5f;
+        r3 = (R0[mad24(height + y, R0Step, x)] - r3) * 0.5f;
+
+        r2 += r4*dy + r6*dx;
+        r3 += r6*dy + r5*dx;
+
+        float scale =
+            c_border[min(x, BORDER_SIZE)] *
+            c_border[min(y, BORDER_SIZE)] *
+            c_border[min(width - x - 1, BORDER_SIZE)] *
+            c_border[min(height - y - 1, BORDER_SIZE)];
+
+        r2 *= scale;
+        r3 *= scale;
+        r4 *= scale;
+        r5 *= scale;
+        r6 *= scale;
+
+        M[mad24(y, mStep, x)] = r4*r4 + r6*r6;
+        M[mad24(height + y, mStep, x)] = (r4 + r5)*r6;
+        M[mad24(2*height + y, mStep, x)] = r5*r5 + r6*r6;
+        M[mad24(3*height + y, mStep, x)] = r4*r2 + r6*r3;
+        M[mad24(4*height + y, mStep, x)] = r6*r2 + r5*r3;
+    }
+}
+
+__kernel void boxFilter5(__global float * dst,
+                         __global const float * src,
+                         __local float * smem,
+                         const int height,  const int width,
+                         int dstStep, int srcStep,
+                         const int ksizeHalf)
+{
+    const int y = get_global_id(1);
+    const int x = get_global_id(0);
+
+    const float boxAreaInv = 1.f / ((1 + 2*ksizeHalf) * (1 + 2*ksizeHalf));
+    const int smw = bdx + 2*ksizeHalf; // shared memory "width"
+    __local float *row = smem + 5 * ty * smw;
+
+    dstStep /= sizeof(*dst);
+    srcStep /= sizeof(*src);
+
+    if (y < height)
+    {
+        // Vertical pass
+        for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx)
+        {
+            int xExt = (int)(bx * bdx) + i - ksizeHalf;
+            xExt = min(max(xExt, 0), width - 1);
+
+#pragma unroll
+            for (int k = 0; k < 5; ++k)
+                row[k*smw + i] = src[mad24(k*height + y, srcStep, xExt)];
+
+            for (int j = 1; j <= ksizeHalf; ++j)
+#pragma unroll
+                for (int k = 0; k < 5; ++k)
+                    row[k*smw + i] +=
+                        src[mad24(k*height + max(y - j, 0), srcStep, xExt)] +
+                        src[mad24(k*height + min(y + j, height - 1), srcStep, xExt)];
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (y < height && y >= 0 && x < width && x >= 0)
+    {
+        // Horizontal pass
+
+        row += tx + ksizeHalf;
+        float res[5];
+
+#pragma unroll
+        for (int k = 0; k < 5; ++k)
+            res[k] = row[k*smw];
+
+        for (int i = 1; i <= ksizeHalf; ++i)
+#pragma unroll
+            for (int k = 0; k < 5; ++k)
+                res[k] += row[k*smw - i] + row[k*smw + i];
+
+#pragma unroll
+        for (int k = 0; k < 5; ++k)
+            dst[mad24(k*height + y, dstStep, x)] = res[k] * boxAreaInv;
+    }
+}
+
+__kernel void updateFlow(__global float4 * flowx, __global float4 * flowy,
+                         __global const float4 * M,
+                         const int height, const int width,
+                         int xStep, int yStep, int mStep)
+{
+    const int y = get_global_id(1);
+    const int x = get_global_id(0);
+
+    xStep /= sizeof(*flowx);
+    yStep /= sizeof(*flowy);
+    mStep /= sizeof(*M);
+
+    if (y < height && y >= 0 && x < width && x >= 0)
+    {
+        float4 g11 = M[mad24(y, mStep, x)];
+        float4 g12 = M[mad24(height + y, mStep, x)];
+        float4 g22 = M[mad24(2*height + y, mStep, x)];
+        float4 h1 =  M[mad24(3*height + y, mStep, x)];
+        float4 h2 =  M[mad24(4*height + y, mStep, x)];
+
+        float4 detInv = (float4)(1.f) / (g11*g22 - g12*g12 + (float4)(1e-3f));
+
+        flowx[mad24(y, xStep, x)] = (g11*h2 - g12*h1) * detInv;
+        flowy[mad24(y, yStep, x)] = (g22*h1 - g12*h2) * detInv;
+    }
+}
+
+__kernel void gaussianBlur5(__global float * dst,
+                            __global const float * src,
+                            __global const float * c_gKer,
+                            __local float * smem,
+                            const int height,  const int width,
+                            int dstStep, int srcStep,
+                            const int ksizeHalf)
+{
+    const int y = get_global_id(1);
+    const int x = get_global_id(0);
+
+    const int smw = bdx + 2*ksizeHalf; // shared memory "width"
+    __local volatile float *row = smem + 5 * ty * smw;
+
+    dstStep /= sizeof(*dst);
+    srcStep /= sizeof(*src);
+
+    if (y < height)
+    {
+        // Vertical pass
+        for (int i = tx; i < bdx + 2*ksizeHalf; i += bdx)
+        {
+            int xExt = (int)(bx * bdx) + i - ksizeHalf;
+            xExt = idx_col(xExt, width - 1);
+
+#pragma unroll
+            for (int k = 0; k < 5; ++k)
+                row[k*smw + i] = src[mad24(k*height + y, srcStep, xExt)] * c_gKer[0];
+
+            for (int j = 1; j <= ksizeHalf; ++j)
+#pragma unroll
+                for (int k = 0; k < 5; ++k)
+                    row[k*smw + i] +=
+                        (src[mad24(k*height + idx_row_low(y - j, height - 1), srcStep, xExt)] +
+                         src[mad24(k*height + idx_row_high(y + j, height - 1), srcStep, xExt)]) * c_gKer[j];
+        }
+    }
+
+    barrier(CLK_LOCAL_MEM_FENCE);
+
+    if (y < height && y >= 0 && x < width && x >= 0)
+    {
+        // Horizontal pass
+
+        row += tx + ksizeHalf;
+        float res[5];
+
+#pragma unroll
+        for (int k = 0; k < 5; ++k)
+            res[k] = row[k*smw] * c_gKer[0];
+
+        for (int i = 1; i <= ksizeHalf; ++i)
+#pragma unroll
+            for (int k = 0; k < 5; ++k)
+                res[k] += (row[k*smw - i] + row[k*smw + i]) * c_gKer[i];
+
+#pragma unroll
+        for (int k = 0; k < 5; ++k)
+            dst[mad24(k*height + y, dstStep, x)] = res[k];
+    }
+}

modules/ocl/src/opencl/stereobm.cl

@@ -258,27 +258,13 @@ float sobel(__global unsigned char *input, int x, int y, int rows, int cols)
 
 float CalcSums(__local float *cols, __local float *cols_cache, int winsz)
 {
-    float cache = 0;
-    float cache2 = 0;
-    int winsz2 = winsz/2;
+    unsigned int cache = cols[0];
 
-    int x = get_local_id(0);
-    int group_size_x = get_local_size(0);
-
-    for(int i = 1; i <= winsz2; i++)
+#pragma unroll
+    for(int i = 1; i <= winsz; i++)
         cache += cols[i];
 
-    cols_cache[0] = cache;
-
-    barrier(CLK_LOCAL_MEM_FENCE);
-
-    if (x < group_size_x - winsz2)
-        cache2 = cols_cache[winsz2];
-    else
-        for(int i = winsz2 + 1; i < winsz; i++)
-            cache2 += cols[i];
-
-    return cols[0] + cache + cache2;
+    return cache;
 }
 
 #define RpT (2 * ROWSperTHREAD)  // got experimentally
@@ -301,8 +287,7 @@ __kernel void textureness_kernel(__global unsigned char *disp, int disp_rows, in
     int beg_row = group_id_y * RpT;
     int end_row = min(beg_row + RpT, disp_rows);
 
-//   if (x < disp_cols)
-//   {
+
     int y = beg_row;
 
     float sum = 0;
@@ -340,11 +325,15 @@ __kernel void textureness_kernel(__global unsigned char *disp, int disp_rows, in
         }
 
         barrier(CLK_LOCAL_MEM_FENCE);
-        float sum_win = CalcSums(cols, cols_cache + local_id_x, winsz) * 255;
-        if (sum_win < threshold)
-            disp[y * disp_step + x] = 0;
+
+        if (x < disp_cols)
+        {
+            float sum_win = CalcSums(cols, cols_cache + local_id_x, winsz) * 255;
+            if (sum_win < threshold)
+                disp[y * disp_step + x] = 0;
+        }
 
         barrier(CLK_LOCAL_MEM_FENCE);
     }
-    //  }
+
 }

modules/ocl/src/opencl/stereobp.cl

@@ -56,8 +56,10 @@
 
 #ifdef T_FLOAT
 #define T float
+#define T4 float4
 #else
 #define T short
+#define T4 short4
 #endif
 
 ///////////////////////////////////////////////////////////////
@@ -71,6 +73,14 @@ T saturate_cast(float v){
 #endif
 }
 
+T4 saturate_cast4(float4 v){
+#ifdef T_SHORT
+    return convert_short4_sat_rte(v); 
+#else
+    return v;
+#endif
+}
+
 #define FLOAT_MAX 3.402823466e+38f
 typedef struct
 {
@@ -84,29 +94,14 @@ typedef struct
 ////////////////////////// comp data //////////////////////////
 ///////////////////////////////////////////////////////////////
 
-float pix_diff_1(__global const uchar *ls, __global const uchar *rs)
+inline float pix_diff_1(const uchar4 l, __global const uchar *rs)
 {
-    return abs((int)(*ls) - *rs); 
+    return abs((int)(l.x) - *rs); 
 }
 
-float pix_diff_3(__global const uchar *ls, __global const uchar *rs)
+float pix_diff_4(const uchar4 l, __global const uchar *rs)
 {
-    const float tr = 0.299f;
-    const float tg = 0.587f;
-    const float tb = 0.114f;
-
-    float val;
-
-    val =  tb * abs((int)ls[0] - rs[0]);
-    val += tg * abs((int)ls[1] - rs[1]);
-    val += tr * abs((int)ls[2] - rs[2]);
-
-    return val;
-}
-float pix_diff_4(__global const uchar *ls, __global const uchar *rs)
-{
-    uchar4 l, r;
-    l = *((__global uchar4 *)ls);
+    uchar4 r;
     r = *((__global uchar4 *)rs);
 
     const float tr = 0.299f;
@@ -122,11 +117,19 @@ float pix_diff_4(__global const uchar *ls, __global const uchar *rs)
     return val;
 }
 
+inline float pix_diff_3(const uchar4 l, __global const uchar *rs)
+{
+    return pix_diff_4(l, rs);
+}
 
 #ifndef CN
 #define CN 4
 #endif
 
+#ifndef CNDISP
+#define CNDISP 64
+#endif
+
 #define CAT(X,Y) X##Y
 #define CAT2(X,Y) CAT(X,Y)
 
@@ -149,19 +152,20 @@ __kernel void comp_data(__global uchar *left,  int left_rows,  int left_cols,  i
         __global T *ds = data + y * data_step + x;
 
         const unsigned int disp_step = data_step * left_rows;
+        const float weightXterm = con_st -> cdata_weight * con_st -> cmax_data_term;
+        const uchar4 ls_data = vload4(0, ls);
 
         for (int disp = 0; disp < con_st -> cndisp; disp++)
         {
             if (x - disp >= 1)
             {
                 float val = 0;
-                val = PIX_DIFF(ls, rs - disp * CN);
-                ds[disp * disp_step] =  saturate_cast(fmin(con_st -> cdata_weight * val, 
-                    con_st -> cdata_weight * con_st -> cmax_data_term));
+                val = PIX_DIFF(ls_data, rs - disp * CN);
+                ds[disp * disp_step] =  saturate_cast(fmin(con_st -> cdata_weight * val, weightXterm));
             }
             else
             {
-                ds[disp * disp_step] =  saturate_cast(con_st -> cdata_weight * con_st -> cmax_data_term);
+                ds[disp * disp_step] =  saturate_cast(weightXterm);
             }
         }
     }
@@ -182,13 +186,20 @@ __kernel void data_step_down(__global T *src, int src_rows,
     {
         src_step /= sizeof(T);
         dst_step /= sizeof(T);
+        int4 coor_step = (int4)(src_rows * src_step);
+        int4 coor = (int4)(min(2*y+0, src_rows-1) * src_step + 2*x+0,
+                           min(2*y+1, src_rows-1) * src_step + 2*x+0,
+                           min(2*y+0, src_rows-1) * src_step + 2*x+1,
+                           min(2*y+1, src_rows-1) * src_step + 2*x+1);
+
         for (int d = 0; d < cndisp; ++d)
         {
             float dst_reg;
-            dst_reg  = src[(d * src_rows + min(2*y+0, src_rows-1)) * src_step + 2*x+0];
-            dst_reg += src[(d * src_rows + min(2*y+1, src_rows-1)) * src_step + 2*x+0];
-            dst_reg += src[(d * src_rows + min(2*y+0, src_rows-1)) * src_step + 2*x+1];
-            dst_reg += src[(d * src_rows + min(2*y+1, src_rows-1)) * src_step + 2*x+1];
+            dst_reg  = src[coor.x];
+            dst_reg += src[coor.y];
+            dst_reg += src[coor.z];
+            dst_reg += src[coor.w];
+            coor += coor_step;
 
             dst[(d * dst_rows + y) * dst_step + x] = saturate_cast(dst_reg);
         }
@@ -224,85 +235,95 @@ __kernel void level_up_message(__global T *src, int src_rows, int src_step,
 ///////////////////////////////////////////////////////////////
 ////////////////////  calc all iterations /////////////////////
 ///////////////////////////////////////////////////////////////
-void calc_min_linear_penalty(__global T * dst, int disp_step, 
-                             int cndisp, float cdisc_single_jump)
+void message(__global T *us_, __global T *ds_, __global T *ls_, __global T *rs_,
+              const __global T *dt,
+              int u_step, int msg_disp_step, int data_disp_step,
+              float4 cmax_disc_term, float4 cdisc_single_jump)
 {
-    float prev = dst[0];
-    float cur;
+    __global T *us = us_ + u_step;
+    __global T *ds = ds_ - u_step;
+    __global T *ls = ls_ + 1;
+    __global T *rs = rs_ - 1;
 
-    for (int disp = 1; disp < cndisp; ++disp)
+    float4 minimum = (float4)(FLOAT_MAX);
+
+    T4 t_dst[CNDISP];
+    float4 dst_reg;
+    float4 prev;
+    float4 cur;
+
+    T t_us = us[0];
+    T t_ds = ds[0];
+    T t_ls = ls[0];
+    T t_rs = rs[0];
+    T t_dt = dt[0];
+
+    prev = (float4)(t_us + t_ls + t_rs + t_dt,
+                    t_ds + t_ls + t_rs + t_dt,
+                    t_us + t_ds + t_rs + t_dt,
+                    t_us + t_ds + t_ls + t_dt);
+
+    minimum = min(prev, minimum);
+
+    t_dst[0] = saturate_cast4(prev);
+
+    for(int i = 1, idx = msg_disp_step; i < CNDISP; ++i, idx+=msg_disp_step)
     {
+        t_us = us[idx];
+        t_ds = ds[idx];
+        t_ls = ls[idx];
+        t_rs = rs[idx];
+        t_dt = dt[data_disp_step * i];
+
+        dst_reg = (float4)(t_us + t_ls + t_rs + t_dt,
+                           t_ds + t_ls + t_rs + t_dt,
+                           t_us + t_ds + t_rs + t_dt,
+                           t_us + t_ds + t_ls + t_dt);
+
+        minimum = min(dst_reg, minimum);
+
         prev += cdisc_single_jump;
-        cur = dst[disp_step * disp];
+        prev = min(prev, dst_reg);
 
-        if (prev < cur)
-        {
-            cur = prev;
-            dst[disp_step * disp] = saturate_cast(prev);
-        }
-
-        prev = cur;
+        t_dst[i] = saturate_cast4(prev);
     }
 
-    prev = dst[(cndisp - 1) * disp_step];
-    for (int disp = cndisp - 2; disp >= 0; disp--)
-    {
-        prev += cdisc_single_jump;
-        cur = dst[disp_step * disp];
-
-        if (prev < cur)
-        {
-            cur = prev;
-            dst[disp_step * disp] = saturate_cast(prev);
-        }
-        prev = cur;
-    }
-}
-void message(const __global T *msg1, const __global T *msg2,
-             const __global T *msg3, const __global T *data, __global T *dst,
-             int msg_disp_step, int data_disp_step, int cndisp, float cmax_disc_term, float cdisc_single_jump)
-{
-    float minimum = FLOAT_MAX;
-
-    for(int i = 0; i < cndisp; ++i)
-    {
-        float dst_reg;
-        dst_reg  = msg1[msg_disp_step * i];
-        dst_reg += msg2[msg_disp_step * i];
-        dst_reg += msg3[msg_disp_step * i];
-        dst_reg += data[data_disp_step * i];
-
-        if (dst_reg < minimum)
-            minimum = dst_reg;
-
-        dst[msg_disp_step * i] = saturate_cast(dst_reg);
-    }
-
-    calc_min_linear_penalty(dst, msg_disp_step, cndisp, cdisc_single_jump);
-
     minimum += cmax_disc_term;
-
-    float sum = 0;
-    for(int i = 0; i < cndisp; ++i)
+    
+    float4 sum = 0;
+    prev = convert_float4(t_dst[CNDISP - 1]);
+    for (int disp = CNDISP - 2; disp >= 0; disp--)
     {
-        float dst_reg = dst[msg_disp_step * i];
-        if (dst_reg > minimum)
-        {
-            dst_reg = minimum;
-            dst[msg_disp_step * i] = saturate_cast(minimum);
-        }
-        sum += dst_reg;
-    }
-    sum /= cndisp;
+        prev += cdisc_single_jump;
+        cur = convert_float4(t_dst[disp]);
+        prev = min(prev, cur);
+        cur = min(prev, minimum);
+        sum += cur;
 
-    for(int i = 0; i < cndisp; ++i)
-        dst[msg_disp_step * i] -= sum;
+        t_dst[disp] = saturate_cast4(cur);
+    }
+
+    dst_reg = convert_float4(t_dst[CNDISP - 1]);
+    dst_reg = min(dst_reg, minimum);
+    t_dst[CNDISP - 1] = saturate_cast4(dst_reg);
+    sum += dst_reg;
+
+    sum /= CNDISP;
+#pragma unroll
+    for(int i = 0, idx = 0; i < CNDISP; ++i, idx+=msg_disp_step)
+    {
+        T4 dst = t_dst[i];
+        us_[idx] = dst.x - sum.x;
+        ds_[idx] = dst.y - sum.y;
+        rs_[idx] = dst.z - sum.z;
+        ls_[idx] = dst.w - sum.w;
+    }
 }
 __kernel void one_iteration(__global T *u,    int u_step,
                             __global T *data, int data_step,
                             __global T *d,    __global T *l, __global T *r,
                             int t, int cols, int rows, 
-                            int cndisp, float cmax_disc_term, float cdisc_single_jump)
+                            float cmax_disc_term, float cdisc_single_jump)
 {
     const int y = get_global_id(1);
     const int x = ((get_global_id(0)) << 1) + ((y + t) & 1);
@@ -321,15 +342,9 @@ __kernel void one_iteration(__global T *u,    int u_step,
         int msg_disp_step = u_step * rows;
         int data_disp_step = data_step * rows;
 
-        message(us + u_step, ls      + 1, rs - 1, dt, us, msg_disp_step, data_disp_step, cndisp, 
-            cmax_disc_term, cdisc_single_jump);
-        message(ds - u_step, ls      + 1, rs - 1, dt, ds, msg_disp_step, data_disp_step, cndisp,
-            cmax_disc_term, cdisc_single_jump);
-
-        message(us + u_step, ds - u_step, rs - 1, dt, rs, msg_disp_step, data_disp_step, cndisp,
-            cmax_disc_term, cdisc_single_jump);
-        message(us + u_step, ds - u_step, ls + 1, dt, ls, msg_disp_step, data_disp_step, cndisp,
-            cmax_disc_term, cdisc_single_jump);
+        message(us, ds, ls, rs, dt,
+                u_step, msg_disp_step, data_disp_step,
+                (float4)(cmax_disc_term), (float4)(cdisc_single_jump));
     }
 }
 

modules/ocl/src/optical_flow_farneback.cpp

@@ -0,0 +1,540 @@
+/*M///////////////////////////////////////////////////////////////////////////////////////
+//
+//  IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
+//
+//  By downloading, copying, installing or using the software you agree to this license.
+//  If you do not agree to this license, do not download, install,
+//  copy or use the software.
+//
+//
+//                           License Agreement
+//                For Open Source Computer Vision Library
+//
+// Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
+// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+//
+// @Authors
+//      Sen Liu, swjtuls1987@126.com
+//
+// Redistribution and use in source and binary forms, with or without modification,
+// are permitted provided that the following conditions are met:
+//
+//   * Redistribution's of source code must retain the above copyright notice,
+//     this list of conditions and the following disclaimer.
+//
+//   * Redistribution's in binary form must reproduce the above copyright notice,
+//     this list of conditions and the following disclaimer in the documentation
+//     and/or other oclMaterials provided with the distribution.
+//
+//   * The name of the copyright holders may not be used to endorse or promote products
+//     derived from this software without specific prior written permission.
+//
+// This software is provided by the copyright holders and contributors "as is" and
+// any express or implied warranties, including, but not limited to, the implied
+// warranties of merchantability and fitness for a particular purpose are disclaimed.
+// In no event shall the Intel Corporation or contributors be liable for any direct,
+// indirect, incidental, special, exemplary, or consequential damages
+// (including, but not limited to, procurement of substitute goods or services;
+// loss of use, data, or profits; or business interruption) however caused
+// and on any theory of liability, whether in contract, strict liability,
+// or tort (including negligence or otherwise) arising in any way out of
+// the use of this software, even if advised of the possibility of such damage.
+//
+//M*/
+
+
+#include "precomp.hpp"
+#include "opencv2/video/tracking.hpp"
+
+using namespace std;
+using namespace cv;
+using namespace cv::ocl;
+
+#define MIN_SIZE 32
+
+namespace cv
+{
+namespace ocl
+{
+///////////////////////////OpenCL kernel strings///////////////////////////
+extern const char *optical_flow_farneback;
+}
+}
+
+namespace cv {
+namespace ocl {
+namespace optflow_farneback
+{
+oclMat g;
+oclMat xg;
+oclMat xxg;
+oclMat gKer;
+
+float ig[4];
+
+inline int divUp(int total, int grain)
+{
+    return (total + grain - 1) / grain;
+}
+
+inline void setGaussianBlurKernel(const float *c_gKer, int ksizeHalf)
+{
+    cv::Mat t_gKer(1, ksizeHalf + 1, CV_32FC1, const_cast<float *>(c_gKer));
+    gKer.upload(t_gKer);
+}
+
+static void gaussianBlurOcl(const oclMat &src, int ksizeHalf, oclMat &dst)
+{
+    string kernelName("gaussianBlur");
+    size_t localThreads[3] = { 256, 1, 1 };
+    size_t globalThreads[3] = { divUp(src.cols, localThreads[0]) * localThreads[0], src.rows, 1 };
+    int smem_size = (localThreads[0] + 2*ksizeHalf) * sizeof(float);
+
+    CV_Assert(dst.size() == src.size());
+    std::vector< std::pair<size_t, const void *> > args;
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&dst.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&src.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&gKer.data));
+    args.push_back(std::make_pair(smem_size, (void *)NULL));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.rows));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.cols));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&ksizeHalf));
+
+    openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName,
+                        globalThreads, localThreads, args, -1, -1);
+}
+
+static void polynomialExpansionOcl(const oclMat &src, int polyN, oclMat &dst)
+{
+    string kernelName("polynomialExpansion");
+    size_t localThreads[3] = { 256, 1, 1 };
+    size_t globalThreads[3] = { divUp(src.cols, localThreads[0] - 2*polyN) * localThreads[0], src.rows, 1 };
+    int smem_size = 3 * localThreads[0] * sizeof(float);
+
+    std::vector< std::pair<size_t, const void *> > args;
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&dst.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&src.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&g.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&xg.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&xxg.data));
+    args.push_back(std::make_pair(smem_size, (void *)NULL));
+    args.push_back(std::make_pair(sizeof(cl_float4), (void *)&ig));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.rows));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.cols));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.step));
+
+    char opt [128];
+    sprintf(opt, "-D polyN=%d", polyN);
+
+    openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName,
+                        globalThreads, localThreads, args, -1, -1, opt);
+}
+
+static void updateMatricesOcl(const oclMat &flowx, const oclMat &flowy, const oclMat &R0, const oclMat &R1, oclMat &M)
+{
+    string kernelName("updateMatrices");
+    size_t localThreads[3] = { 32, 8, 1 };
+    size_t globalThreads[3] = { divUp(flowx.cols, localThreads[0]) * localThreads[0],
+                                divUp(flowx.rows, localThreads[1]) * localThreads[1],
+                                1
+                              };
+
+    std::vector< std::pair<size_t, const void *> > args;
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&M.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&flowx.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&flowy.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&R0.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&R1.data));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowx.rows));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowx.cols));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&M.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowx.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowy.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&R0.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&R1.step));
+
+    openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName,
+                        globalThreads, localThreads, args, -1, -1);
+}
+
+static void boxFilter5Ocl(const oclMat &src, int ksizeHalf, oclMat &dst)
+{
+    string kernelName("boxFilter5");
+    int height = src.rows / 5;
+    size_t localThreads[3] = { 256, 1, 1 };
+    size_t globalThreads[3] = { divUp(src.cols, localThreads[0]) * localThreads[0], height, 1 };
+    int smem_size = (localThreads[0] + 2*ksizeHalf) * 5 * sizeof(float);
+
+    std::vector< std::pair<size_t, const void *> > args;
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&dst.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&src.data));
+    args.push_back(std::make_pair(smem_size, (void *)NULL));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&height));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.cols));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&ksizeHalf));
+
+    openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName,
+                        globalThreads, localThreads, args, -1, -1);
+}
+
+static void updateFlowOcl(const oclMat &M, oclMat &flowx, oclMat &flowy)
+{
+    string kernelName("updateFlow");
+    int cols = divUp(flowx.cols, 4);
+    size_t localThreads[3] = { 32, 8, 1 };
+    size_t globalThreads[3] = { divUp(cols, localThreads[0]) * localThreads[0],
+                                divUp(flowx.rows, localThreads[1]) * localThreads[0],
+                                1
+                              };
+
+    std::vector< std::pair<size_t, const void *> > args;
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&flowx.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&flowy.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&M.data));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowx.rows));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&cols));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowx.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&flowy.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&M.step));
+
+    openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName,
+                        globalThreads, localThreads, args, -1, -1);
+}
+
+static void gaussianBlur5Ocl(const oclMat &src, int ksizeHalf, oclMat &dst)
+{
+    string kernelName("gaussianBlur5");
+    int height = src.rows / 5;
+    int width = src.cols;
+    size_t localThreads[3] = { 256, 1, 1 };
+    size_t globalThreads[3] = { divUp(width, localThreads[0]) * localThreads[0], height, 1 };
+    int smem_size = (localThreads[0] + 2*ksizeHalf) * 5 * sizeof(float);
+
+    std::vector< std::pair<size_t, const void *> > args;
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&dst.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&src.data));
+    args.push_back(std::make_pair(sizeof(cl_mem), (void *)&gKer.data));
+    args.push_back(std::make_pair(smem_size, (void *)NULL));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&height));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&width));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&dst.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&src.step));
+    args.push_back(std::make_pair(sizeof(cl_int), (void *)&ksizeHalf));
+
+    openCLExecuteKernel(Context::getContext(), &optical_flow_farneback, kernelName,
+                        globalThreads, localThreads, args, -1, -1);
+}
+}
+}
+} // namespace cv { namespace ocl { namespace optflow_farneback
+
+static oclMat allocMatFromBuf(int rows, int cols, int type, oclMat &mat)
+{
+    if (!mat.empty() && mat.type() == type && mat.rows >= rows && mat.cols >= cols)
+        return mat(Rect(0, 0, cols, rows));
+    return mat = oclMat(rows, cols, type);
+}
+
+cv::ocl::FarnebackOpticalFlow::FarnebackOpticalFlow()
+{
+    numLevels = 5;
+    pyrScale = 0.5;
+    fastPyramids = false;
+    winSize = 13;
+    numIters = 10;
+    polyN = 5;
+    polySigma = 1.1;
+    flags = 0;
+}
+
+void cv::ocl::FarnebackOpticalFlow::releaseMemory()
+{
+    frames_[0].release();
+    frames_[1].release();
+    pyrLevel_[0].release();
+    pyrLevel_[1].release();
+    M_.release();
+    bufM_.release();
+    R_[0].release();
+    R_[1].release();
+    blurredFrame_[0].release();
+    blurredFrame_[1].release();
+    pyramid0_.clear();
+    pyramid1_.clear();
+}
+
+void cv::ocl::FarnebackOpticalFlow::prepareGaussian(
+    int n, double sigma, float *g, float *xg, float *xxg,
+    double &ig11, double &ig03, double &ig33, double &ig55)
+{
+    double s = 0.;
+    for (int x = -n; x <= n; x++)
+    {
+        g[x] = (float)std::exp(-x*x/(2*sigma*sigma));
+        s += g[x];
+    }
+
+    s = 1./s;
+    for (int x = -n; x <= n; x++)
+    {
+        g[x] = (float)(g[x]*s);
+        xg[x] = (float)(x*g[x]);
+        xxg[x] = (float)(x*x*g[x]);
+    }
+
+    Mat_<double> G(6, 6);
+    G.setTo(0);
+
+    for (int y = -n; y <= n; y++)
+    {
+        for (int x = -n; x <= n; x++)
+        {
+            G(0,0) += g[y]*g[x];
+            G(1,1) += g[y]*g[x]*x*x;
+            G(3,3) += g[y]*g[x]*x*x*x*x;
+            G(5,5) += g[y]*g[x]*x*x*y*y;
+        }
+    }
+
+    //G[0][0] = 1.;
+    G(2,2) = G(0,3) = G(0,4) = G(3,0) = G(4,0) = G(1,1);
+    G(4,4) = G(3,3);
+    G(3,4) = G(4,3) = G(5,5);
+
+    // invG:
+    // [ x        e  e    ]
+    // [    y             ]
+    // [       y          ]
+    // [ e        z       ]
+    // [ e           z    ]
+    // [                u ]
+    Mat_<double> invG = G.inv(DECOMP_CHOLESKY);
+
+    ig11 = invG(1,1);
+    ig03 = invG(0,3);
+    ig33 = invG(3,3);
+    ig55 = invG(5,5);
+}
+
+void cv::ocl::FarnebackOpticalFlow::setPolynomialExpansionConsts(int n, double sigma)
+{
+    vector<float> buf(n*6 + 3);
+    float* g = &buf[0] + n;
+    float* xg = g + n*2 + 1;
+    float* xxg = xg + n*2 + 1;
+
+    if (sigma < FLT_EPSILON)
+        sigma = n*0.3;
+
+    double ig11, ig03, ig33, ig55;
+    prepareGaussian(n, sigma, g, xg, xxg, ig11, ig03, ig33, ig55);
+
+    cv::Mat t_g(1, n + 1, CV_32FC1, g);
+    cv::Mat t_xg(1, n + 1, CV_32FC1, xg);
+    cv::Mat t_xxg(1, n + 1, CV_32FC1, xxg);
+
+    optflow_farneback::g.upload(t_g);
+    optflow_farneback::xg.upload(t_xg);
+    optflow_farneback::xxg.upload(t_xxg);
+
+    optflow_farneback::ig[0] = static_cast<float>(ig11);
+    optflow_farneback::ig[1] = static_cast<float>(ig03);
+    optflow_farneback::ig[2] = static_cast<float>(ig33);
+    optflow_farneback::ig[3] = static_cast<float>(ig55);
+}
+
+void cv::ocl::FarnebackOpticalFlow::updateFlow_boxFilter(
+    const oclMat& R0, const oclMat& R1, oclMat& flowx, oclMat &flowy,
+    oclMat& M, oclMat &bufM, int blockSize, bool updateMatrices)
+{
+    optflow_farneback::boxFilter5Ocl(M, blockSize/2, bufM);
+
+    swap(M, bufM);
+
+    finish();
+
+    optflow_farneback::updateFlowOcl(M, flowx, flowy);
+
+    if (updateMatrices)
+        optflow_farneback::updateMatricesOcl(flowx, flowy, R0, R1, M);
+}
+
+
+void cv::ocl::FarnebackOpticalFlow::updateFlow_gaussianBlur(
+    const oclMat& R0, const oclMat& R1, oclMat& flowx, oclMat& flowy,
+    oclMat& M, oclMat &bufM, int blockSize, bool updateMatrices)
+{
+    optflow_farneback::gaussianBlur5Ocl(M, blockSize/2, bufM);
+
+    swap(M, bufM);
+
+    optflow_farneback::updateFlowOcl(M, flowx, flowy);
+
+    if (updateMatrices)
+        optflow_farneback::updateMatricesOcl(flowx, flowy, R0, R1, M);
+}
+
+
+void cv::ocl::FarnebackOpticalFlow::operator ()(
+    const oclMat &frame0, const oclMat &frame1, oclMat &flowx, oclMat &flowy)
+{
+    CV_Assert(frame0.channels() == 1 && frame1.channels() == 1);
+    CV_Assert(frame0.size() == frame1.size());
+    CV_Assert(polyN == 5 || polyN == 7);
+    CV_Assert(!fastPyramids || std::abs(pyrScale - 0.5) < 1e-6);
+
+    Size size = frame0.size();
+    oclMat prevFlowX, prevFlowY, curFlowX, curFlowY;
+
+    flowx.create(size, CV_32F);
+    flowy.create(size, CV_32F);
+    oclMat flowx0 = flowx;
+    oclMat flowy0 = flowy;
+
+    // Crop unnecessary levels
+    double scale = 1;
+    int numLevelsCropped = 0;
+    for (; numLevelsCropped < numLevels; numLevelsCropped++)
+    {
+        scale *= pyrScale;
+        if (size.width*scale < MIN_SIZE || size.height*scale < MIN_SIZE)
+            break;
+    }
+
+    frame0.convertTo(frames_[0], CV_32F);
+    frame1.convertTo(frames_[1], CV_32F);
+
+    if (fastPyramids)
+    {
+        // Build Gaussian pyramids using pyrDown()
+        pyramid0_.resize(numLevelsCropped + 1);
+        pyramid1_.resize(numLevelsCropped + 1);
+        pyramid0_[0] = frames_[0];
+        pyramid1_[0] = frames_[1];
+        for (int i = 1; i <= numLevelsCropped; ++i)
+        {
+            pyrDown(pyramid0_[i - 1], pyramid0_[i]);
+            pyrDown(pyramid1_[i - 1], pyramid1_[i]);
+        }
+    }
+
+    setPolynomialExpansionConsts(polyN, polySigma);
+
+    for (int k = numLevelsCropped; k >= 0; k--)
+    {
+        scale = 1;
+        for (int i = 0; i < k; i++)
+            scale *= pyrScale;
+
+        double sigma = (1./scale - 1) * 0.5;
+        int smoothSize = cvRound(sigma*5) | 1;
+        smoothSize = std::max(smoothSize, 3);
+
+        int width = cvRound(size.width*scale);
+        int height = cvRound(size.height*scale);
+
+        if (fastPyramids)
+        {
+            width = pyramid0_[k].cols;
+            height = pyramid0_[k].rows;
+        }
+
+        if (k > 0)
+        {
+            curFlowX.create(height, width, CV_32F);
+            curFlowY.create(height, width, CV_32F);
+        }
+        else
+        {
+            curFlowX = flowx0;
+            curFlowY = flowy0;
+        }
+
+        if (!prevFlowX.data)
+        {
+            if (flags & cv::OPTFLOW_USE_INITIAL_FLOW)
+            {
+                resize(flowx0, curFlowX, Size(width, height), 0, 0, INTER_LINEAR);
+                resize(flowy0, curFlowY, Size(width, height), 0, 0, INTER_LINEAR);
+                multiply(scale, curFlowX, curFlowX);
+                multiply(scale, curFlowY, curFlowY);
+            }
+            else
+            {
+                curFlowX.setTo(0);
+                curFlowY.setTo(0);
+            }
+        }
+        else
+        {
+            resize(prevFlowX, curFlowX, Size(width, height), 0, 0, INTER_LINEAR);
+            resize(prevFlowY, curFlowY, Size(width, height), 0, 0, INTER_LINEAR);
+            multiply(1./pyrScale, curFlowX, curFlowX);
+            multiply(1./pyrScale, curFlowY, curFlowY);
+        }
+
+        oclMat M = allocMatFromBuf(5*height, width, CV_32F, M_);
+        oclMat bufM = allocMatFromBuf(5*height, width, CV_32F, bufM_);
+        oclMat R[2] =
+        {
+            allocMatFromBuf(5*height, width, CV_32F, R_[0]),
+            allocMatFromBuf(5*height, width, CV_32F, R_[1])
+        };
+
+        if (fastPyramids)
+        {
+            optflow_farneback::polynomialExpansionOcl(pyramid0_[k], polyN, R[0]);
+            optflow_farneback::polynomialExpansionOcl(pyramid1_[k], polyN, R[1]);
+        }
+        else
+        {
+            oclMat blurredFrame[2] =
+            {
+                allocMatFromBuf(size.height, size.width, CV_32F, blurredFrame_[0]),
+                allocMatFromBuf(size.height, size.width, CV_32F, blurredFrame_[1])
+            };
+            oclMat pyrLevel[2] =
+            {
+                allocMatFromBuf(height, width, CV_32F, pyrLevel_[0]),
+                allocMatFromBuf(height, width, CV_32F, pyrLevel_[1])
+            };
+
+            Mat g = getGaussianKernel(smoothSize, sigma, CV_32F);
+            optflow_farneback::setGaussianBlurKernel(g.ptr<float>(smoothSize/2), smoothSize/2);
+
+            for (int i = 0; i < 2; i++)
+            {
+                optflow_farneback::gaussianBlurOcl(frames_[i], smoothSize/2, blurredFrame[i]);
+                resize(blurredFrame[i], pyrLevel[i], Size(width, height), INTER_LINEAR);
+                optflow_farneback::polynomialExpansionOcl(pyrLevel[i], polyN, R[i]);
+            }
+        }
+
+        optflow_farneback::updateMatricesOcl(curFlowX, curFlowY, R[0], R[1], M);
+
+        if (flags & OPTFLOW_FARNEBACK_GAUSSIAN)
+        {
+            Mat g = getGaussianKernel(winSize, winSize/2*0.3f, CV_32F);
+            optflow_farneback::setGaussianBlurKernel(g.ptr<float>(winSize/2), winSize/2);
+        }
+        for (int i = 0; i < numIters; i++)
+        {
+            if (flags & OPTFLOW_FARNEBACK_GAUSSIAN)
+                updateFlow_gaussianBlur(R[0], R[1], curFlowX, curFlowY, M, bufM, winSize, i < numIters-1);
+            else
+                updateFlow_boxFilter(R[0], R[1], curFlowX, curFlowY, M, bufM, winSize, i < numIters-1);
+        }
+
+        prevFlowX = curFlowX;
+        prevFlowY = curFlowY;
+    }
+
+    flowx = curFlowX;
+    flowy = curFlowY;
+}

modules/ocl/src/stereobp.cpp

@@ -236,13 +236,13 @@ namespace cv
                 args.push_back( make_pair( sizeof(cl_int) , (void *)&t));
                 args.push_back( make_pair( sizeof(cl_int) , (void *)&cols));
                 args.push_back( make_pair( sizeof(cl_int) , (void *)&rows));
-                args.push_back( make_pair( sizeof(cl_int) , (void *)&cndisp));
                 args.push_back( make_pair( sizeof(cl_float) , (void *)&cmax_disc_term));
                 args.push_back( make_pair( sizeof(cl_float) , (void *)&cdisc_single_jump));
 
                 size_t gt[3] = {cols, rows, 1}, lt[3] = {16, 16, 1};
-                const char* t_opt  = data_type == CV_16S ? "-D T_SHORT":"-D T_FLOAT";
-                openCLExecuteKernel(clCxt, &stereobp, kernelName, gt, lt, args, -1, -1, t_opt);
+                char opt[80] = "";
+                sprintf(opt, "-D %s -D CNDISP=%d", data_type == CV_16S ? "T_SHORT":"T_FLOAT", cndisp);
+                openCLExecuteKernel(clCxt, &stereobp, kernelName, gt, lt, args, -1, -1, opt);
             }
 
             static void calc_all_iterations_calls(int cols, int rows, int iters, oclMat &u,

modules/ocl/test/main.cpp

@@ -73,14 +73,12 @@ void print_info()
 #endif
 
 }
-std::string workdir;
 int main(int argc, char **argv)
 {
-    TS::ptr()->init("ocl");
+    TS::ptr()->init(".");
     InitGoogleTest(&argc, argv);
     const char *keys =
         "{ h | help     | false              | print help message }"
-        "{ w | workdir  | ../../../samples/c/| set working directory }"
         "{ t | type     | gpu                | set device type:cpu or gpu}"
         "{ p | platform | 0                  | set platform id }"
         "{ d | device   | 0                  | set device id }";
@@ -92,7 +90,6 @@ int main(int argc, char **argv)
         cmd.printParams();
         return 0;
     }
-    workdir = cmd.get<string>("workdir");
     string type = cmd.get<string>("type");
     unsigned int pid = cmd.get<unsigned int>("platform");
     int device = cmd.get<int>("device");
@@ -118,6 +115,8 @@ int main(int argc, char **argv)
 
     setDevice(oclinfo[pid], device);
 
+    setBinaryDiskCache(CACHE_UPDATE);
+
     cout << "Device type:" << type << endl << "Device name:" << oclinfo[pid].DeviceName[device] << endl;
     return RUN_ALL_TESTS();
 }

modules/ocl/test/test_calib3d.cpp

@@ -50,7 +50,6 @@
 
 using namespace cv;
 
-extern std::string workdir;
 PARAM_TEST_CASE(StereoMatchBM, int, int)
 {
     int n_disp;
@@ -66,9 +65,9 @@ PARAM_TEST_CASE(StereoMatchBM, int, int)
 TEST_P(StereoMatchBM, Regression)
 {
 
-    Mat left_image  = readImage("stereobm/aloe-L.png", IMREAD_GRAYSCALE);
-    Mat right_image = readImage("stereobm/aloe-R.png", IMREAD_GRAYSCALE);
-    Mat disp_gold   = readImage("stereobm/aloe-disp.png", IMREAD_GRAYSCALE);
+    Mat left_image  = readImage("gpu/stereobm/aloe-L.png", IMREAD_GRAYSCALE);
+    Mat right_image = readImage("gpu/stereobm/aloe-R.png", IMREAD_GRAYSCALE);
+    Mat disp_gold   = readImage("gpu/stereobm/aloe-disp.png", IMREAD_GRAYSCALE);
     ocl::oclMat d_left, d_right;
     ocl::oclMat d_disp(left_image.size(), CV_8U);
     Mat  disp;
@@ -113,9 +112,9 @@ PARAM_TEST_CASE(StereoMatchBP, int, int, int, float, float, float, float)
 };
 TEST_P(StereoMatchBP, Regression)
 {
-    Mat left_image  = readImage("stereobp/aloe-L.png");
-    Mat right_image = readImage("stereobp/aloe-R.png");
-    Mat disp_gold   = readImage("stereobp/aloe-disp.png", IMREAD_GRAYSCALE);
+    Mat left_image  = readImage("gpu/stereobp/aloe-L.png");
+    Mat right_image = readImage("gpu/stereobp/aloe-R.png");
+    Mat disp_gold   = readImage("gpu/stereobp/aloe-disp.png", IMREAD_GRAYSCALE);
     ocl::oclMat d_left, d_right;
     ocl::oclMat d_disp;
     Mat  disp;
@@ -166,9 +165,9 @@ PARAM_TEST_CASE(StereoMatchConstSpaceBP, int, int, int, int, float, float, float
 };
 TEST_P(StereoMatchConstSpaceBP, Regression)
 {
-    Mat left_image  = readImage("csstereobp/aloe-L.png");
-    Mat right_image = readImage("csstereobp/aloe-R.png");
-    Mat disp_gold   = readImage("csstereobp/aloe-disp.png", IMREAD_GRAYSCALE);
+    Mat left_image  = readImage("gpu/csstereobp/aloe-L.png");
+    Mat right_image = readImage("gpu/csstereobp/aloe-R.png");
+    Mat disp_gold   = readImage("gpu/csstereobp/aloe-disp.png", IMREAD_GRAYSCALE);
 
     ocl::oclMat d_left, d_right;
     ocl::oclMat d_disp;

modules/ocl/test/test_canny.cpp

@@ -48,7 +48,6 @@
 
 ////////////////////////////////////////////////////////
 // Canny
-extern std::string workdir;
 IMPLEMENT_PARAM_CLASS(AppertureSize, int);
 IMPLEMENT_PARAM_CLASS(L2gradient, bool);
 
@@ -67,7 +66,7 @@ PARAM_TEST_CASE(Canny, AppertureSize, L2gradient)
 
 TEST_P(Canny, Accuracy)
 {
-    cv::Mat img = readImage(workdir + "fruits.jpg", cv::IMREAD_GRAYSCALE);
+    cv::Mat img = readImage("cv/shared/fruits.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(img.empty());
 
     double low_thresh = 50.0;

modules/ocl/test/test_moments.cpp

@@ -45,7 +45,7 @@ TEST_P(MomentsTest, Mat)
     {
         if(test_contours)
         {
-            Mat src = imread( workdir + "../cpp/pic3.png", IMREAD_GRAYSCALE );
+            Mat src = readImage( "cv/shared/pic3.png", IMREAD_GRAYSCALE );
             ASSERT_FALSE(src.empty());
             Mat canny_output;
             vector<vector<Point> > contours;

modules/ocl/test/test_objdetect.cpp

@@ -63,11 +63,8 @@ PARAM_TEST_CASE(HOG, Size, int)
     {
         winSize = GET_PARAM(0);
         type = GET_PARAM(1);
-        img_rgb = readImage(workdir + "../gpu/road.png");
-        if(img_rgb.empty())
-        {
-            std::cout << "Couldn't read road.png" << std::endl;
-        }
+        img_rgb = readImage("gpu/hog/road.png");
+        ASSERT_FALSE(img_rgb.empty());
     }
 };
 
@@ -146,17 +143,17 @@ TEST_P(HOG, Detect)
     if (winSize.width == 48 && winSize.height == 96)
     {
         // daimler's base
-        ocl_hog.setSVMDetector(ocl_hog.getPeopleDetector48x96());
+        ocl_hog.setSVMDetector(hog.getDaimlerPeopleDetector());
         hog.setSVMDetector(hog.getDaimlerPeopleDetector());
     }
     else if (winSize.width == 64 && winSize.height == 128)
     {
-        ocl_hog.setSVMDetector(ocl_hog.getPeopleDetector64x128());
+        ocl_hog.setSVMDetector(hog.getDefaultPeopleDetector());
         hog.setSVMDetector(hog.getDefaultPeopleDetector());
     }
     else
     {
-        ocl_hog.setSVMDetector(ocl_hog.getDefaultPeopleDetector());
+        ocl_hog.setSVMDetector(hog.getDefaultPeopleDetector());
         hog.setSVMDetector(hog.getDefaultPeopleDetector());
     }
 
@@ -211,18 +208,11 @@ PARAM_TEST_CASE(Haar, int, CascadeName)
     virtual void SetUp()
     {
         flags = GET_PARAM(0);
-        cascadeName = (workdir + "../../data/haarcascades/").append(GET_PARAM(1));
-        if( (!cascade.load( cascadeName )) || (!cpucascade.load(cascadeName)) )
-        {
-            std::cout << "ERROR: Could not load classifier cascade" << std::endl;
-            return;
-        }
-        img = readImage(workdir + "lena.jpg", IMREAD_GRAYSCALE);
-        if(img.empty())
-        {
-            std::cout << "Couldn't read lena.jpg" << std::endl;
-            return ;
-        }
+        cascadeName = (string(cvtest::TS::ptr()->get_data_path()) + "cv/cascadeandhog/cascades/").append(GET_PARAM(1));
+        ASSERT_TRUE(cascade.load( cascadeName ));
+        ASSERT_TRUE(cpucascade.load(cascadeName));
+        img = readImage("cv/shared/lena.png", IMREAD_GRAYSCALE);
+        ASSERT_FALSE(img.empty());
         equalizeHist(img, img);
         d_img.upload(img);
     }

modules/ocl/test/test_optflow.cpp

@@ -75,7 +75,7 @@ PARAM_TEST_CASE(GoodFeaturesToTrack, MinDistance)
 
 TEST_P(GoodFeaturesToTrack, Accuracy)
 {
-    cv::Mat frame = readImage(workdir + "../gpu/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
+    cv::Mat frame = readImage("gpu/opticalflow/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(frame.empty());
 
     int maxCorners = 1000;
@@ -146,10 +146,10 @@ PARAM_TEST_CASE(TVL1, bool)
 
 TEST_P(TVL1, Accuracy)
 {
-    cv::Mat frame0 = readImage(workdir + "../gpu/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
+    cv::Mat frame0 = readImage("gpu/opticalflow/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(frame0.empty());
 
-    cv::Mat frame1 = readImage(workdir + "../gpu/rubberwhale2.png", cv::IMREAD_GRAYSCALE);
+    cv::Mat frame1 = readImage("gpu/opticalflow/rubberwhale2.png", cv::IMREAD_GRAYSCALE);
     ASSERT_FALSE(frame1.empty());
 
     cv::ocl::OpticalFlowDual_TVL1_OCL d_alg;
@@ -188,10 +188,10 @@ PARAM_TEST_CASE(Sparse, bool, bool)
 
 TEST_P(Sparse, Mat)
 {
-    cv::Mat frame0 = readImage(workdir + "../gpu/rubberwhale1.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
+    cv::Mat frame0 = readImage("gpu/opticalflow/rubberwhale1.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
     ASSERT_FALSE(frame0.empty());
 
-    cv::Mat frame1 = readImage(workdir + "../gpu/rubberwhale2.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
+    cv::Mat frame1 = readImage("gpu/opticalflow/rubberwhale2.png", useGray ? cv::IMREAD_GRAYSCALE : cv::IMREAD_COLOR);
     ASSERT_FALSE(frame1.empty());
 
     cv::Mat gray_frame;
@@ -272,6 +272,78 @@ TEST_P(Sparse, Mat)
 INSTANTIATE_TEST_CASE_P(OCL_Video, Sparse, Combine(
     Values(false, true),
     Values(false, true)));
+//////////////////////////////////////////////////////
+// FarnebackOpticalFlow
+
+namespace
+{
+    IMPLEMENT_PARAM_CLASS(PyrScale, double)
+        IMPLEMENT_PARAM_CLASS(PolyN, int)
+        CV_FLAGS(FarnebackOptFlowFlags, 0, OPTFLOW_FARNEBACK_GAUSSIAN)
+        IMPLEMENT_PARAM_CLASS(UseInitFlow, bool)
+}
+
+PARAM_TEST_CASE(Farneback, PyrScale, PolyN, FarnebackOptFlowFlags, UseInitFlow)
+{
+    double pyrScale;
+    int polyN;
+    int flags;
+    bool useInitFlow;
+
+    virtual void SetUp()
+    {
+        pyrScale = GET_PARAM(0);
+        polyN = GET_PARAM(1);
+        flags = GET_PARAM(2);
+        useInitFlow = GET_PARAM(3);
+    }
+};
+
+TEST_P(Farneback, Accuracy)
+{
+    cv::Mat frame0 = readImage("gpu/opticalflow/rubberwhale1.png", cv::IMREAD_GRAYSCALE);
+    ASSERT_FALSE(frame0.empty());
+
+    cv::Mat frame1 = readImage("gpu/opticalflow/rubberwhale2.png", cv::IMREAD_GRAYSCALE);
+    ASSERT_FALSE(frame1.empty());
+
+    double polySigma = polyN <= 5 ? 1.1 : 1.5;
+
+    cv::ocl::FarnebackOpticalFlow farn;
+    farn.pyrScale = pyrScale;
+    farn.polyN = polyN;
+    farn.polySigma = polySigma;
+    farn.flags = flags;
+
+    cv::ocl::oclMat d_flowx, d_flowy;
+    farn(oclMat(frame0), oclMat(frame1), d_flowx, d_flowy);
+
+    cv::Mat flow;
+    if (useInitFlow)
+    {
+        cv::Mat flowxy[] = {cv::Mat(d_flowx), cv::Mat(d_flowy)};
+        cv::merge(flowxy, 2, flow);
+
+        farn.flags |= cv::OPTFLOW_USE_INITIAL_FLOW;
+        farn(oclMat(frame0), oclMat(frame1), d_flowx, d_flowy);
+    }
+
+    cv::calcOpticalFlowFarneback(
+        frame0, frame1, flow, farn.pyrScale, farn.numLevels, farn.winSize,
+        farn.numIters, farn.polyN, farn.polySigma, farn.flags);
+
+    std::vector<cv::Mat> flowxy;
+    cv::split(flow, flowxy);
+
+    EXPECT_MAT_SIMILAR(flowxy[0], d_flowx, 0.1);
+    EXPECT_MAT_SIMILAR(flowxy[1], d_flowy, 0.1);
+}
+
+INSTANTIATE_TEST_CASE_P(OCL_Video, Farneback, testing::Combine(
+    testing::Values(PyrScale(0.3), PyrScale(0.5), PyrScale(0.8)),
+    testing::Values(PolyN(5), PolyN(7)),
+    testing::Values(FarnebackOptFlowFlags(0), FarnebackOptFlowFlags(cv::OPTFLOW_FARNEBACK_GAUSSIAN)),
+    testing::Values(UseInitFlow(false), UseInitFlow(true))));
 
 #endif // HAVE_OPENCL
 

modules/ts/misc/xls-report.py

@@ -40,9 +40,15 @@
            Corollary 2: an empty 'properties' dictionary matches every property set.
 
         3) If a matching matcher is found, its 'name' string is presumed to be the name
-           of the configuration the XML file corresponds to. Otherwise, a warning is
-           printed. A warning is also printed if two different property sets match to the
-           same configuration name.
+           of the configuration the XML file corresponds to. A warning is printed if
+           two different property sets match to the same configuration name.
+
+        4) If a such a matcher isn't found, if --include-unmatched was specified, the
+           configuration name is assumed to be the relative path from the sheet's
+           directory to the XML file's containing directory. If the XML file isinstance
+           directly inside the sheet's directory, the configuration name is instead
+           a dump of all its properties. If --include-unmatched wasn't specified,
+           the XML file is ignored and a warning is printed.
 
     * 'configurations': [string]
         List of names for compile-time and runtime configurations of OpenCV.
@@ -67,6 +73,7 @@
 from __future__ import division
 
 import ast
+import errno
 import fnmatch
 import logging
 import numbers
@@ -74,7 +81,6 @@ import os, os.path
 import re
 
 from argparse import ArgumentParser
-from collections import OrderedDict
 from glob import glob
 from itertools import ifilter
 
@@ -96,18 +102,25 @@ error_speedup_style = xlwt.easyxf('pattern: pattern solid, fore_color orange')
 header_style = xlwt.easyxf('font: bold true; alignment: horizontal centre, vertical top, wrap True')
 
 class Collector(object):
-    def __init__(self, config_match_func):
+    def __init__(self, config_match_func, include_unmatched):
         self.__config_cache = {}
         self.config_match_func = config_match_func
+        self.include_unmatched = include_unmatched
         self.tests = {}
+        self.extra_configurations = set()
 
     # Format a sorted sequence of pairs as if it was a dictionary.
     # We can't just use a dictionary instead, since we want to preserve the sorted order of the keys.
     @staticmethod
-    def __format_config_cache_key(pairs):
-        return '{' + ', '.join(repr(k) + ': ' + repr(v) for (k, v) in pairs) + '}'
+    def __format_config_cache_key(pairs, multiline=False):
+        return (
+          ('{\n' if multiline else '{') +
+          (',\n' if multiline else ', ').join(
+             ('  ' if multiline else '') + repr(k) + ': ' + repr(v) for (k, v) in pairs) +
+          ('\n}\n' if multiline else '}')
+        )
 
-    def collect_from(self, xml_path):
+    def collect_from(self, xml_path, default_configuration):
         run = parseLogFile(xml_path)
 
         module = run.properties['module_name']
@@ -123,8 +136,17 @@ class Collector(object):
             configuration = self.config_match_func(properties)
 
             if configuration is None:
-                logging.warning('failed to match properties to a configuration: %s',
-                    Collector.__format_config_cache_key(props_key))
+                if self.include_unmatched:
+                    if default_configuration is not None:
+                        configuration = default_configuration
+                    else:
+                        configuration = Collector.__format_config_cache_key(props_key, multiline=True)
+
+                    self.extra_configurations.add(configuration)
+                else:
+                    logging.warning('failed to match properties to a configuration: %s',
+                        Collector.__format_config_cache_key(props_key))
+
             else:
                 same_config_props = [it[0] for it in self.__config_cache.iteritems() if it[1] == configuration]
                 if len(same_config_props) > 0:
@@ -137,11 +159,17 @@ class Collector(object):
 
         if configuration is None: return
 
-        module_tests = self.tests.setdefault(module, OrderedDict())
+        module_tests = self.tests.setdefault(module, {})
 
         for test in run.tests:
             test_results = module_tests.setdefault((test.shortName(), test.param()), {})
-            test_results[configuration] = test.get("gmean") if test.status == 'run' else test.status
+            new_result = test.get("gmean") if test.status == 'run' else test.status
+            test_results[configuration] = min(
+              test_results.get(configuration), new_result,
+              key=lambda r: (1, r) if isinstance(r, numbers.Number) else
+                            (2,) if r is not None else
+                            (3,)
+            ) # prefer lower result; prefer numbers to errors and errors to nothing
 
 def make_match_func(matchers):
     def match_func(properties):
@@ -159,6 +187,8 @@ def main():
     arg_parser.add_argument('sheet_dirs', nargs='+', metavar='DIR', help='directory containing perf test logs')
     arg_parser.add_argument('-o', '--output', metavar='XLS', default='report.xls', help='name of output file')
     arg_parser.add_argument('-c', '--config', metavar='CONF', help='global configuration file')
+    arg_parser.add_argument('--include-unmatched', action='store_true',
+        help='include results from XML files that were not recognized by configuration matchers')
 
     args = arg_parser.parse_args()
 
@@ -176,7 +206,8 @@ def main():
         try:
             with open(os.path.join(sheet_path, 'sheet.conf')) as sheet_conf_file:
                 sheet_conf = ast.literal_eval(sheet_conf_file.read())
-        except Exception:
+        except IOError as ioe:
+            if ioe.errno != errno.ENOENT: raise
             sheet_conf = {}
             logging.debug('no sheet.conf for %s', sheet_path)
 
@@ -185,12 +216,18 @@ def main():
         config_names = sheet_conf.get('configurations', [])
         config_matchers = sheet_conf.get('configuration_matchers', [])
 
-        collector = Collector(make_match_func(config_matchers))
+        collector = Collector(make_match_func(config_matchers), args.include_unmatched)
 
         for root, _, filenames in os.walk(sheet_path):
             logging.info('looking in %s', root)
             for filename in fnmatch.filter(filenames, '*.xml'):
-                collector.collect_from(os.path.join(root, filename))
+                if os.path.normpath(sheet_path) == os.path.normpath(root):
+                  default_conf = None
+                else:
+                  default_conf = os.path.relpath(root, sheet_path)
+                collector.collect_from(os.path.join(root, filename), default_conf)
+
+        config_names.extend(sorted(collector.extra_configurations - set(config_names)))
 
         sheet = wb.add_sheet(sheet_conf.get('sheet_name', os.path.basename(os.path.abspath(sheet_path))))
 
@@ -203,7 +240,7 @@ def main():
         sheet_comparisons = sheet_conf.get('comparisons', [])
 
         for i, w in enumerate([2000, 15000, 2500, 2000, 15000]
-                + (len(config_names) + 1 + len(sheet_comparisons)) * [3000]):
+                + (len(config_names) + 1 + len(sheet_comparisons)) * [4000]):
             sheet.col(i).width = w
 
         for i, caption in enumerate(['Module', 'Test', 'Image\nsize', 'Data\ntype', 'Parameters']
@@ -218,7 +255,7 @@ def main():
                          for module, color in module_colors.iteritems()}
 
         for module, tests in sorted(collector.tests.iteritems()):
-            for ((test, param), configs) in tests.iteritems():
+            for ((test, param), configs) in sorted(tests.iteritems()):
                 sheet.write(row, 0, module, module_styles.get(module, xlwt.Style.default_style))
                 sheet.write(row, 1, test)
 

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="28@ANDROID_PLATFORM_VERSION_CODE@"
-    android:versionName="2.8" >
+    android:versionCode="29@ANDROID_PLATFORM_VERSION_CODE@"
+    android:versionName="2.9" >
 
     <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.8_<platform>.apk
+    adb install OpenCV-2.4.6-android-sdk/apk/OpenCV_2.4.6_Manager_2.9_<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.8_armv7a-neon.apk          |
+| 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.8_armv7a-neon-android8.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.8_armeabi.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.8_x86.apk                  |
+| Intel x86                    |    >= 2.3    | OpenCV_2.4.6_Manager_2.9_x86.apk                  |
 +------------------------------+--------------+---------------------------------------------------+
-| MIPS                         |    >= 2.3    | OpenCV_2.4.6_Manager_2.8_mips.apk                 |
+| MIPS                         |    >= 2.3    | OpenCV_2.4.6_Manager_2.9_mips.apk                 |
 +------------------------------+--------------+---------------------------------------------------+

platforms/scripts/ABI_compat_generator.py

@@ -6,7 +6,7 @@ import os
 
 
 architecture = 'armeabi'
-excludedHeaders = set(['hdf5.h', 'cap_ios.h', 'eigen.hpp', 'cxeigen.hpp']) #TOREMOVE
+excludedHeaders = set(['hdf5.h', 'cap_ios.h', 'ios.h', 'eigen.hpp', 'cxeigen.hpp']) #TOREMOVE
 systemIncludes = ['sources/cxx-stl/gnu-libstdc++/4.6/include', \
     '/opt/android-ndk-r8c/platforms/android-8/arch-arm', # TODO: check if this one could be passed as command line arg
     'sources/cxx-stl/gnu-libstdc++/4.6/libs/armeabi-v7a/include']

samples/android/native-activity/.cproject

@@ -20,7 +20,7 @@
 					<folderInfo id="0.882924228." name="/" resourcePath="">
 						<toolChain id="org.eclipse.cdt.build.core.prefbase.toolchain.1667980868" name="No ToolChain" resourceTypeBasedDiscovery="false" superClass="org.eclipse.cdt.build.core.prefbase.toolchain">
 							<targetPlatform id="org.eclipse.cdt.build.core.prefbase.toolchain.1667980868.2108168132" name=""/>
-							<builder autoBuildTarget="" command="&quot;${NDKROOT}/ndk-build.cmd&quot;" enableAutoBuild="true" enableCleanBuild="false" id="org.eclipse.cdt.build.core.settings.default.builder.328915772" incrementalBuildTarget="" keepEnvironmentInBuildfile="false" managedBuildOn="false" name="Gnu Make Builder" superClass="org.eclipse.cdt.build.core.settings.default.builder"/>
+							<builder autoBuildTarget="" command="${NDKROOT}/ndk-build.cmd" enableAutoBuild="true" enableCleanBuild="false" id="org.eclipse.cdt.build.core.settings.default.builder.328915772" incrementalBuildTarget="" keepEnvironmentInBuildfile="false" managedBuildOn="false" name="Gnu Make Builder" superClass="org.eclipse.cdt.build.core.settings.default.builder"/>
 							<tool id="org.eclipse.cdt.build.core.settings.holder.libs.630148311" name="holder for library settings" superClass="org.eclipse.cdt.build.core.settings.holder.libs"/>
 							<tool id="org.eclipse.cdt.build.core.settings.holder.525090327" name="Assembly" superClass="org.eclipse.cdt.build.core.settings.holder">
 								<inputType id="org.eclipse.cdt.build.core.settings.holder.inType.1491216279" languageId="org.eclipse.cdt.core.assembly" languageName="Assembly" sourceContentType="org.eclipse.cdt.core.asmSource" superClass="org.eclipse.cdt.build.core.settings.holder.inType"/>

samples/android/native-activity/AndroidManifest.xml

@@ -10,6 +10,7 @@
 
         <activity android:name="CvNativeActivity"
                   android:label="@string/app_name"
+                  android:screenOrientation="landscape"
                   android:configChanges="orientation|keyboardHidden">
             <intent-filter>
                 <action android:name="android.intent.action.MAIN" />
@@ -17,7 +18,9 @@
                 </intent-filter>
             </activity>
         <activity android:name="android.app.NativeActivity"
-                  android:label="@string/app_name">
+                  android:label="@string/app_name"
+                  android:screenOrientation="landscape"
+                  android:configChanges="keyboardHidden|orientation">
             <meta-data android:name="android.app.lib_name"
                     android:value="native_activity" />
         </activity>

samples/android/native-activity/jni/Android.mk

@@ -7,7 +7,7 @@ include ../../sdk/native/jni/OpenCV.mk
 LOCAL_MODULE    := native_activity
 LOCAL_SRC_FILES := native.cpp
 LOCAL_LDLIBS    += -lm -llog -landroid
-LOCAL_STATIC_LIBRARIES := android_native_app_glue
+LOCAL_STATIC_LIBRARIES += android_native_app_glue
 
 include $(BUILD_SHARED_LIBRARY)
 

samples/android/native-activity/jni/native.cpp

@@ -78,18 +78,29 @@ static void engine_draw_frame(Engine* engine, const cv::Mat& frame)
         return;
     }
 
-    void* pixels = buffer.bits;
+    int32_t* pixels = (int32_t*)buffer.bits;
 
     int left_indent = (buffer.width-frame.cols)/2;
     int top_indent = (buffer.height-frame.rows)/2;
 
-    for (int yy = top_indent; yy < std::min(frame.rows+top_indent, buffer.height); yy++)
+    if (top_indent > 0)
     {
-        unsigned char* line = (unsigned char*)pixels + left_indent*4*sizeof(unsigned char);
-        size_t line_size = std::min(frame.cols, buffer.width)*4*sizeof(unsigned char);
+        memset(pixels, 0, top_indent*buffer.stride*sizeof(int32_t));
+        pixels += top_indent*buffer.stride;
+    }
+
+    for (int yy = 0; yy < frame.rows; yy++)
+    {
+        if (left_indent > 0)
+        {
+            memset(pixels, 0, left_indent*sizeof(int32_t));
+            memset(pixels+left_indent+frame.cols, 0, (buffer.stride-frame.cols-left_indent)*sizeof(int32_t));
+        }
+        int32_t* line = pixels + left_indent;
+        size_t line_size = frame.cols*4*sizeof(unsigned char);
         memcpy(line, frame.ptr<unsigned char>(yy), line_size);
         // go to next line
-        pixels = (int32_t*)pixels + buffer.stride;
+        pixels += buffer.stride;
     }
     ANativeWindow_unlockAndPost(engine->app->window);
 }

samples/ocl/CMakeLists.txt

@@ -27,7 +27,7 @@ if(BUILD_EXAMPLES AND OCV_DEPENDENCIES_FOUND)
     target_link_libraries(${the_target} ${OPENCV_LINKER_LIBS} ${OPENCV_OCL_SAMPLES_REQUIRED_DEPS})
 
     set_target_properties(${the_target} PROPERTIES
-      OUTPUT_NAME "${name}_${project}"
+      OUTPUT_NAME "${project}-example-${name}"
       PROJECT_LABEL "(EXAMPLE_${project_upper}) ${name}")
 
     if(ENABLE_SOLUTION_FOLDERS)

samples/ocl/clahe.cpp

@@ -0,0 +1,108 @@
+#include <iostream>
+#include "opencv2/core/core.hpp"
+#include "opencv2/imgproc/imgproc.hpp"
+#include "opencv2/highgui/highgui.hpp"
+#include "opencv2/ocl/ocl.hpp"
+using namespace cv;
+using namespace std;
+
+Ptr<CLAHE> pFilter;
+int tilesize;
+int cliplimit;
+string outfile;
+
+static void TSize_Callback(int pos)
+{
+    if(pos==0)
+    {
+        pFilter->setTilesGridSize(Size(1,1));
+    }
+    pFilter->setTilesGridSize(Size(tilesize,tilesize));
+}
+
+static void Clip_Callback(int)
+{
+    pFilter->setClipLimit(cliplimit);
+}
+
+int main(int argc, char** argv)
+{
+    const char* keys =
+        "{ i | input   |                    | specify input image }"
+        "{ c | camera  |    0               | specify camera id   }"
+        "{ s | use_cpu |    false           | use cpu algorithm   }"
+        "{ o | output  | clahe_output.jpg   | specify output save path}";
+
+    CommandLineParser cmd(argc, argv, keys);
+    string infile = cmd.get<string>("i");
+    outfile = cmd.get<string>("o");
+    int camid = cmd.get<int>("c");
+    bool use_cpu = cmd.get<bool>("s");
+    CvCapture* capture = 0;
+    bool running = true;
+
+    namedWindow("CLAHE");
+    createTrackbar("Tile Size", "CLAHE", &tilesize, 32, (TrackbarCallback)TSize_Callback);
+    createTrackbar("Clip Limit", "CLAHE", &cliplimit, 20, (TrackbarCallback)Clip_Callback);
+    Mat frame, outframe;
+    ocl::oclMat d_outframe;
+    
+    int cur_clip;
+    Size cur_tilesize;
+    if(use_cpu)
+    {
+        pFilter = createCLAHE();
+    }
+    else
+    {
+        pFilter = ocl::createCLAHE();
+    }
+    cur_clip = (int)pFilter->getClipLimit();
+    cur_tilesize = pFilter->getTilesGridSize();
+    setTrackbarPos("Tile Size", "CLAHE", cur_tilesize.width);
+    setTrackbarPos("Clip Limit", "CLAHE", cur_clip);
+    if(infile != "")
+    {
+        frame = imread(infile);
+        if(frame.empty())
+        {
+            cout << "error read image: " << infile << endl;
+            return -1;
+        }
+    }
+    else
+    {
+        capture = cvCaptureFromCAM(camid);
+    }
+    cout << "\nControls:\n"
+         << "\to - save output image\n"
+         << "\tESC - exit\n";
+    while(running)
+    {
+        if(capture)
+            frame = cvQueryFrame(capture);
+        else
+            frame = imread(infile);
+        if(frame.empty())
+        {
+            continue;
+        }
+        if(use_cpu)
+        {
+            cvtColor(frame, frame, COLOR_BGR2GRAY);
+            pFilter->apply(frame, outframe);
+        }
+        else
+        {
+            ocl::oclMat d_frame(frame);
+            ocl::cvtColor(d_frame, d_outframe, COLOR_BGR2GRAY);
+            pFilter->apply(d_outframe, d_outframe);
+            d_outframe.download(outframe);
+        }
+        imshow("CLAHE", outframe);
+        char key = (char)cvWaitKey(3);
+        if(key == 'o') imwrite(outfile, outframe);
+        else if(key == 27) running = false;
+    }
+    return 0;
+}

samples/ocl/facedetect.cpp

@@ -252,8 +252,13 @@ void Draw(Mat& img, vector<Rect>& faces, double scale)
         radius = cvRound((r->width + r->height)*0.25*scale);
         circle( img, center, radius, color, 3, 8, 0 );
     }
-    imshow( "result", img );
     imwrite( outputName, img );
+    if(abs(scale-1.0)>.001)
+    {
+        resize(img, img, Size((int)(img.cols/scale), (int)(img.rows/scale)));
+    }
+    imshow( "result", img );
+    
 }
 
 

samples/ocl/hog.cpp

@@ -57,6 +57,7 @@ private:
     string vdo_source;
     string output;
     int camera_id;
+    bool write_once;
 };
 
 int main(int argc, char** argv)
@@ -97,6 +98,7 @@ App::App(CommandLineParser& cmd)
          << "\tESC - exit\n"
          << "\tm - change mode GPU <-> CPU\n"
          << "\tg - convert image to gray or not\n"
+         << "\to - save output image once, or switch on/off video save\n"
          << "\t1/q - increase/decrease HOG scale\n"
          << "\t2/w - increase/decrease levels count\n"
          << "\t3/e - increase/decrease HOG group threshold\n"
@@ -120,6 +122,7 @@ App::App(CommandLineParser& cmd)
     hit_threshold = win_width == 48 ? 1.4 : 0.;
     scale = 1.05;
     gamma_corr = true;
+    write_once = false;
 
     cout << "Group threshold: " << gr_threshold << endl;
     cout << "Levels number: " << nlevels << endl;
@@ -254,10 +257,11 @@ void App::run()
 
             workEnd();
 
-            if (output!="")
+            if (output!="" && write_once)
             {
                 if (img_source!="")     // wirte image
                 {
+                    write_once = false;
                     imwrite(output, img_to_show);
                 }
                 else                    //write video
@@ -340,6 +344,10 @@ void App::handleKey(char key)
         gamma_corr = !gamma_corr;
         cout << "Gamma correction: " << gamma_corr << endl;
         break;
+    case 'o':
+    case 'O':
+        write_once = !write_once;
+        break;
     }
 }
 

samples/ocl/stereo_match.cpp

@@ -49,7 +49,7 @@ struct App
         return ss.str();
     }
 private:
-    bool running;
+    bool running, write_once;
 
     Mat left_src, right_src;
     Mat left, right;
@@ -115,6 +115,7 @@ App::App(CommandLineParser& cmd)
     cout << "stereo_match_ocl sample\n";
     cout << "\nControls:\n"
          << "\tesc - exit\n"
+         << "\to - save output image once\n"
          << "\tp - print current parameters\n"
          << "\tg - convert source images into gray\n"
          << "\tm - change stereo match method\n"
@@ -132,6 +133,7 @@ App::App(CommandLineParser& cmd)
     else cout << "unknown method!\n";
     ndisp = cmd.get<int>("n");
     out_img = cmd.get<string>("o");
+    write_once = false;
 }
 
 
@@ -161,10 +163,8 @@ void App::run()
     printParams();
 
     running = true;
-    bool written = false;
     while (running)
     {
-
         // Prepare disparity map of specified type
         Mat disp;
         oclMat d_disp;
@@ -192,19 +192,21 @@ void App::run()
             csbp(d_left, d_right, d_disp);
             break;
         }
+
         // Show results
         d_disp.download(disp);
         workEnd();
+
         if (method != BM)
         {
             disp.convertTo(disp, 0);
         }
         putText(disp, text(), Point(5, 25), FONT_HERSHEY_SIMPLEX, 1.0, Scalar::all(255));
         imshow("disparity", disp);
-        if(!written)
+        if(write_once)
         {
             imwrite(out_img, disp);
-            written = true;
+            write_once = false;
         }
         handleKey((char)waitKey(3));
     }
@@ -378,6 +380,10 @@ void App::handleKey(char key)
             cout << "level_count: " << csbp.levels << endl;
         }
         break;
+    case 'o':
+    case 'O':
+        write_once = true;
+        break;
     }
 }