OpenCV Manager version++ for 2.6.4.2r2 update.

(cherry picked from commit aa00b5a0ab)

Conflicts:

	platforms/android/service/engine/jni/BinderComponent/TegraDetector.cpp

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/androidcamera/camera_wrapper/camera_wrapper.cpp

@@ -1,5 +1,8 @@
-#if !defined(ANDROID_r2_2_0) && !defined(ANDROID_r2_3_3) && !defined(ANDROID_r3_0_1) && !defined(ANDROID_r4_0_0) && !defined(ANDROID_r4_0_3) && !defined(ANDROID_r4_1_1) && !defined(ANDROID_r4_2_0)
-# error Building camera wrapper for your version of Android is not supported by OpenCV. You need to modify OpenCV sources in order to compile camera wrapper for your version of Android.
+#if !defined(ANDROID_r2_2_0) && !defined(ANDROID_r2_3_3) && !defined(ANDROID_r3_0_1) && \
+ !defined(ANDROID_r4_0_0) && !defined(ANDROID_r4_0_3) && !defined(ANDROID_r4_1_1) && \
+ !defined(ANDROID_r4_2_0) && !defined(ANDROID_r4_3_0)
+# error Building camera wrapper for your version of Android is not supported by OpenCV.\
+ You need to modify OpenCV sources in order to compile camera wrapper for your version of Android.
 #endif
 
 #include <camera/Camera.h>
@@ -16,17 +19,18 @@
 //Include SurfaceTexture.h file with the SurfaceTexture class
 # include <gui/SurfaceTexture.h>
 # define MAGIC_OPENCV_TEXTURE_ID (0x10)
-#else // defined(ANDROID_r3_0_1) || defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3)
-//TODO: This is either 2.2 or 2.3. Include the headers for ISurface.h access
-#if defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0)
-#include <gui/ISurface.h>
-#include <gui/BufferQueue.h>
+#elif defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0)
+# include <gui/ISurface.h>
+# include <gui/BufferQueue.h>
+#elif defined(ANDROID_r4_3_0)
+# include <gui/IGraphicBufferProducer.h>
+# include <gui/BufferQueue.h>
 #else
 # include <surfaceflinger/ISurface.h>
-#endif  // defined(ANDROID_r4_1_1)
-#endif  // defined(ANDROID_r3_0_1) || defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3)
+#endif
 
 #include <string>
+#include <fstream>
 
 //undef logging macro from /system/core/libcutils/loghack.h
 #ifdef LOGD
@@ -45,7 +49,6 @@
 # undef LOGE
 #endif
 
-
 // LOGGING
 #include <android/log.h>
 #define CAMERA_LOG_TAG "OpenCV_NativeCamera"
@@ -60,7 +63,7 @@ using namespace android;
 
 void debugShowFPS();
 
-#if defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0)
+#if defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0)
 class ConsumerListenerStub: public BufferQueue::ConsumerListener
 {
 public:
@@ -73,6 +76,29 @@ public:
 };
 #endif
 
+std::string getProcessName()
+{
+    std::string result;
+    std::ifstream f;
+
+    f.open("/proc/self/cmdline");
+    if (f.is_open())
+    {
+        std::string fullPath;
+        std::getline(f, fullPath, '\0');
+        if (!fullPath.empty())
+        {
+            int i = fullPath.size()-1;
+            while ((i >= 0) && (fullPath[i] != '/')) i--;
+            result = fullPath.substr(i+1, std::string::npos);
+        }
+    }
+
+    f.close();
+
+    return result;
+}
+
 void debugShowFPS()
 {
     static int mFrameCount = 0;
@@ -280,7 +306,7 @@ public:
     }
 
     virtual void postData(int32_t msgType, const sp<IMemory>& dataPtr
-    #if defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3) || defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0)
+    #if defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3) || defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0)
                           ,camera_frame_metadata_t*
 #endif
                           )
@@ -361,7 +387,9 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
     typedef sp<Camera> (*Android22ConnectFuncType)();
     typedef sp<Camera> (*Android23ConnectFuncType)(int);
     typedef sp<Camera> (*Android3DConnectFuncType)(int, int);
+    typedef sp<Camera> (*Android43ConnectFuncType)(int, const String16&, int);
 
+    const int ANY_CAMERA_INDEX = -1;
     const int BACK_CAMERA_INDEX = 99;
     const int FRONT_CAMERA_INDEX = 98;
 
@@ -372,14 +400,24 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
     CAMERA_SUPPORT_MODE_ZSL = 0x08 /* Camera Sensor supports ZSL mode. */
     };
 
+    // used for Android 4.3
+    enum {
+        USE_CALLING_UID = -1
+    };
+
     const char Android22ConnectName[] = "_ZN7android6Camera7connectEv";
     const char Android23ConnectName[] = "_ZN7android6Camera7connectEi";
     const char Android3DConnectName[] = "_ZN7android6Camera7connectEii";
+    const char Android43ConnectName[] = "_ZN7android6Camera7connectEiRKNS_8String16Ei";
 
     int localCameraIndex = cameraId;
 
+    if (cameraId == ANY_CAMERA_INDEX)
+    {
+        localCameraIndex = 0;
+    }
 #if !defined(ANDROID_r2_2_0)
-    if (cameraId == BACK_CAMERA_INDEX)
+    else if (cameraId == BACK_CAMERA_INDEX)
     {
         LOGD("Back camera selected");
         for (int i = 0; i < Camera::getNumberOfCameras(); i++)
@@ -450,6 +488,12 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
         LOGD("Connecting to CameraService v 3D");
         camera = Android3DConnect(localCameraIndex, CAMERA_SUPPORT_MODE_2D);
     }
+    else if (Android43ConnectFuncType Android43Connect = (Android43ConnectFuncType)dlsym(CameraHALHandle, Android43ConnectName))
+    {
+        std::string currentProcName = getProcessName();
+        LOGD("Current process name for camera init: %s", currentProcName.c_str());
+        camera = Android43Connect(localCameraIndex, String16(currentProcName.c_str()), USE_CALLING_UID);
+    }
     else
     {
         dlclose(CameraHALHandle);
@@ -471,7 +515,7 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
     handler->camera = camera;
     handler->cameraId = localCameraIndex;
 
-    if (prevCameraParameters != 0)
+    if (prevCameraParameters != NULL)
     {
         LOGI("initCameraConnect: Setting paramers from previous camera handler");
         camera->setParameters(prevCameraParameters->flatten());
@@ -503,11 +547,11 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
         const char* available_focus_modes = handler->params.get(CameraParameters::KEY_SUPPORTED_FOCUS_MODES);
         if (available_focus_modes != 0)
         {
-        if (strstr(available_focus_modes, "continuous-video") != NULL)
-        {
-        handler->params.set(CameraParameters::KEY_FOCUS_MODE, CameraParameters::FOCUS_MODE_CONTINUOUS_VIDEO);
+            if (strstr(available_focus_modes, "continuous-video") != NULL)
+            {
+                handler->params.set(CameraParameters::KEY_FOCUS_MODE, CameraParameters::FOCUS_MODE_CONTINUOUS_VIDEO);
 
-        status_t resParams = handler->camera->setParameters(handler->params.flatten());
+                status_t resParams = handler->camera->setParameters(handler->params.flatten());
 
                 if (resParams != 0)
                 {
@@ -517,8 +561,8 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
                 {
                     LOGD("initCameraConnect: autofocus is set to mode \"continuous-video\"");
                 }
+            }
         }
-    }
 #endif
 
         //check if yuv420sp format available. Set this format as preview format.
@@ -560,26 +604,25 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
         }
     }
 
-    status_t pdstatus;
+    status_t bufferStatus;
 #if defined(ANDROID_r2_2_0)
-    pdstatus = camera->setPreviewDisplay(sp<ISurface>(0 /*new DummySurface*/));
-    if (pdstatus != 0)
-        LOGE("initCameraConnect: failed setPreviewDisplay(0) call; camera migth not work correctly on some devices");
+    bufferStatus = camera->setPreviewDisplay(sp<ISurface>(0 /*new DummySurface*/));
+    if (bufferStatus != 0)
+        LOGE("initCameraConnect: failed setPreviewDisplay(0) call (status %d); camera might not work correctly on some devices", bufferStatus);
 #elif defined(ANDROID_r2_3_3)
     /* Do nothing in case of 2.3 for now */
-
 #elif defined(ANDROID_r3_0_1) || defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3)
     sp<SurfaceTexture> surfaceTexture = new SurfaceTexture(MAGIC_OPENCV_TEXTURE_ID);
-    pdstatus = camera->setPreviewTexture(surfaceTexture);
-    if (pdstatus != 0)
-        LOGE("initCameraConnect: failed setPreviewTexture call; camera migth not work correctly");
-#elif defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0)
+    bufferStatus = camera->setPreviewTexture(surfaceTexture);
+    if (bufferStatus != 0)
+        LOGE("initCameraConnect: failed setPreviewTexture call (status %d); camera might not work correctly", bufferStatus);
+#elif defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0)
     sp<BufferQueue> bufferQueue = new BufferQueue();
     sp<BufferQueue::ConsumerListener> queueListener = new ConsumerListenerStub();
     bufferQueue->consumerConnect(queueListener);
-    pdstatus = camera->setPreviewTexture(bufferQueue);
-    if (pdstatus != 0)
-    LOGE("initCameraConnect: failed setPreviewTexture call; camera migth not work correctly");
+    bufferStatus = camera->setPreviewTexture(bufferQueue);
+    if (bufferStatus != 0)
+        LOGE("initCameraConnect: failed setPreviewTexture call; camera might not work correctly");
 #endif
 
 #if (defined(ANDROID_r2_2_0) || defined(ANDROID_r2_3_3) || defined(ANDROID_r3_0_1))
@@ -595,9 +638,9 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
 #endif //!(defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3))
 
     LOGD("Starting preview");
-    status_t resStart = camera->startPreview();
+    status_t previewStatus = camera->startPreview();
 
-    if (resStart != 0)
+    if (previewStatus != 0)
     {
         LOGE("initCameraConnect: startPreview() fails. Closing camera connection...");
         handler->closeCameraConnect();
@@ -605,7 +648,7 @@ CameraHandler* CameraHandler::initCameraConnect(const CameraCallback& callback,
     }
     else
     {
-    LOGD("Preview started successfully");
+        LOGD("Preview started successfully");
     }
 
     return handler;
@@ -620,9 +663,11 @@ void CameraHandler::closeCameraConnect()
     }
 
     camera->stopPreview();
+#if defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3) || defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0)
+    camera->setPreviewCallbackFlags(CAMERA_FRAME_CALLBACK_FLAG_NOOP);
+#endif
     camera->disconnect();
     camera.clear();
-
     camera=NULL;
     // ATTENTION!!!!!!!!!!!!!!!!!!!!!!!!!!
     // When we set
@@ -863,14 +908,60 @@ void CameraHandler::applyProperties(CameraHandler** ppcameraHandler)
 
     if (*ppcameraHandler == 0)
     {
-        LOGE("applyProperties: Passed null *ppcameraHandler");
+        LOGE("applyProperties: Passed NULL *ppcameraHandler");
         return;
     }
 
-    LOGD("CameraHandler::applyProperties()");
-    CameraHandler* previousCameraHandler=*ppcameraHandler;
-    CameraParameters curCameraParameters(previousCameraHandler->params.flatten());
+    CameraParameters curCameraParameters((*ppcameraHandler)->params.flatten());
 
+#if defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3) || defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0)
+    CameraHandler* handler=*ppcameraHandler;
+
+    handler->camera->stopPreview();
+    handler->camera->setPreviewCallbackFlags(CAMERA_FRAME_CALLBACK_FLAG_NOOP);
+
+    status_t reconnectStatus = handler->camera->reconnect();
+    if (reconnectStatus != 0)
+    {
+        LOGE("applyProperties: failed to reconnect camera (status %d)", reconnectStatus);
+        return;
+    }
+
+    handler->camera->setParameters(curCameraParameters.flatten());
+    handler->params.unflatten(curCameraParameters.flatten());
+
+    status_t bufferStatus;
+# if defined(ANDROID_r4_0_0) || defined(ANDROID_r4_0_3)
+    sp<SurfaceTexture> surfaceTexture = new SurfaceTexture(MAGIC_OPENCV_TEXTURE_ID);
+    bufferStatus = handler->camera->setPreviewTexture(surfaceTexture);
+    if (bufferStatus != 0)
+        LOGE("applyProperties: failed setPreviewTexture call (status %d); camera might not work correctly", bufferStatus);
+# elif defined(ANDROID_r4_1_1) || defined(ANDROID_r4_2_0) || defined(ANDROID_r4_3_0)
+    sp<BufferQueue> bufferQueue = new BufferQueue();
+    sp<BufferQueue::ConsumerListener> queueListener = new ConsumerListenerStub();
+    bufferQueue->consumerConnect(queueListener);
+    bufferStatus = handler->camera->setPreviewTexture(bufferQueue);
+    if (bufferStatus != 0)
+        LOGE("applyProperties: failed setPreviewTexture call; camera might not work correctly");
+# endif
+
+    handler->camera->setPreviewCallbackFlags( CAMERA_FRAME_CALLBACK_FLAG_ENABLE_MASK | CAMERA_FRAME_CALLBACK_FLAG_COPY_OUT_MASK);//with copy
+
+    LOGD("Starting preview");
+    status_t previewStatus = handler->camera->startPreview();
+
+    if (previewStatus != 0)
+    {
+        LOGE("initCameraConnect: startPreview() fails. Closing camera connection...");
+        handler->closeCameraConnect();
+        handler = NULL;
+    }
+    else
+    {
+        LOGD("Preview started successfully");
+    }
+#else
+    CameraHandler* previousCameraHandler=*ppcameraHandler;
     CameraCallback cameraCallback=previousCameraHandler->cameraCallback;
     void* userData=previousCameraHandler->userData;
     int cameraId=previousCameraHandler->cameraId;
@@ -879,7 +970,6 @@ void CameraHandler::applyProperties(CameraHandler** ppcameraHandler)
     previousCameraHandler->closeCameraConnect();
     LOGD("CameraHandler::applyProperties(): after previousCameraHandler->closeCameraConnect");
 
-
     LOGD("CameraHandler::applyProperties(): before initCameraConnect");
     CameraHandler* handler=initCameraConnect(cameraCallback, cameraId, userData, &curCameraParameters);
     LOGD("CameraHandler::applyProperties(): after initCameraConnect, handler=0x%x", (int)handler);
@@ -892,6 +982,7 @@ void CameraHandler::applyProperties(CameraHandler** ppcameraHandler)
         }
     }
     (*ppcameraHandler)=handler;
+#endif
 }
 
 

modules/androidcamera/src/camera_activity.cpp

@@ -431,14 +431,14 @@ void CameraActivity::applyProperties()
 int CameraActivity::getFrameWidth()
 {
     if (frameWidth <= 0)
-    frameWidth = getProperty(ANDROID_CAMERA_PROPERTY_FRAMEWIDTH);
+        frameWidth = getProperty(ANDROID_CAMERA_PROPERTY_FRAMEWIDTH);
     return frameWidth;
 }
 
 int CameraActivity::getFrameHeight()
 {
     if (frameHeight <= 0)
-    frameHeight = getProperty(ANDROID_CAMERA_PROPERTY_FRAMEHEIGHT);
+        frameHeight = getProperty(ANDROID_CAMERA_PROPERTY_FRAMEHEIGHT);
     return frameHeight;
 }
 

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 2
 
 #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="211@ANDROID_PLATFORM_VERSION_CODE@"
+    android:versionName="2.11" >
 
     <uses-sdk android:minSdkVersion="@ANDROID_NATIVE_API_LEVEL@" />
     <uses-feature android:name="android.hardware.touchscreen" android:required="false"/>

platforms/android/service/engine/CMakeLists.txt

@@ -26,19 +26,32 @@ endif()
 
 configure_file("${CMAKE_CURRENT_SOURCE_DIR}/${ANDROID_MANIFEST_FILE}" "${OpenCV_BINARY_DIR}/platforms/android/service/engine/.build/${ANDROID_MANIFEST_FILE}"  @ONLY)
 
-link_directories("${ANDROID_SOURCE_TREE}/out/target/product/generic/system/lib" "${ANDROID_SOURCE_TREE}/out/target/product/${ANDROID_PRODUCT}/system/lib" "${ANDROID_SOURCE_TREE}/bin/${ANDROID_ARCH_NAME}")
+link_directories(
+  "${ANDROID_SOURCE_TREE}/out/target/product/generic/system/lib"
+  "${ANDROID_SOURCE_TREE}/out/target/product/${ANDROID_PRODUCT}/system/lib"
+  "${ANDROID_SOURCE_TREE}/bin/${ANDROID_ARCH_NAME}")
+
+file(GLOB engine_files "jni/BinderComponent/*.cpp" "jni/BinderComponent/*.h" "jni/include/*.h")
+set(engine_libs "z" "binder" "log" "utils")
+
+if (TEGRA_DETECTOR)
+  if (ANDROID_NATIVE_API_LEVEL GREATER 8)
+    add_definitions(-DUSE_TEGRA_HW_DETECTOR)
+    list(APPEND engine_libs ${TEGRA_DETECTOR} GLESv2 EGL)
+  else()
+    message(FATAL_ERROR "Tegra detector required native api level 9 or above")
+  endif()
+endif()
 
 # -D__SUPPORT_ARMEABI_FEATURES key is also available
 add_definitions(-DPLATFORM_ANDROID -D__SUPPORT_ARMEABI_V7A_FEATURES -D__SUPPORT_TEGRA3 -D__SUPPORT_MIPS)
 set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fno-rtti -fno-exceptions")
-
 set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} -Wl,-allow-shlib-undefined")
-file(GLOB engine_files "jni/BinderComponent/*.cpp" "jni/BinderComponent/*.h" "jni/include/*.h")
 
-include_directories(jni/BinderComponent jni/include)
+include_directories("jni/BinderComponent" "jni/include")
 include_directories(SYSTEM "${ANDROID_SOURCE_TREE}/frameworks/base/include" "${ANDROID_SOURCE_TREE}/system/core/include")
 add_library(${engine} SHARED ${engine_files})
-target_link_libraries(${engine} z binder log utils)
+target_link_libraries(${engine} ${engine_libs})
 
 set_target_properties(${engine} PROPERTIES
     OUTPUT_NAME ${engine}
@@ -51,7 +64,15 @@ add_custom_command(TARGET ${engine} POST_BUILD COMMAND ${CMAKE_STRIP} --strip-un
 file(GLOB engine_jni_files "jni/JNIWrapper/*.cpp" "jni/JNIWrapper/*.h" "jni/include/*.h")
 list(APPEND engine_jni_files jni/NativeService/CommonPackageManager.cpp jni/NativeService/PackageInfo.cpp)
 
-include_directories(jni/include jni/JNIWrapper jni/NativeService jni/BinderComponent "${ANDROID_SOURCE_TREE}/frameworks/base/include" "${ANDROID_SOURCE_TREE}/system/core/include" "${ANDROID_SOURCE_TREE}/frameworks/base/core/jni")
+include_directories(
+  jni/include jni/JNIWrapper
+  jni/NativeService
+  jni/BinderComponent
+  "${ANDROID_SOURCE_TREE}/frameworks/base/include"
+  "${ANDROID_SOURCE_TREE}/system/core/include"
+  "${ANDROID_SOURCE_TREE}/frameworks/base/core/jni"
+  )
+
 add_library(${engine}_jni SHARED ${engine_jni_files})
 target_link_libraries(${engine}_jni z binder log utils android_runtime ${engine})
 

platforms/android/service/engine/jni/BinderComponent/HardwareDetector.cpp

@@ -13,7 +13,7 @@ int GetCpuID()
     map<string, string> cpu_info = GetCpuInfo();
     map<string, string>::const_iterator it;
 
-    #if defined(__i386__)
+#if defined(__i386__)
     LOGD("Using X86 HW detector");
     result |= ARCH_X86;
     it = cpu_info.find("flags");
@@ -161,8 +161,11 @@ int GetProcessorCount()
 
 int DetectKnownPlatforms()
 {
+#if defined(__arm__) && defined(USE_TEGRA_HW_DETECTOR)
     int tegra_status = DetectTegra();
-
+#else
+    int tegra_status = NOT_TEGRA;
+#endif
     // All Tegra platforms since Tegra3
     if (2 < tegra_status)
     {
@@ -172,4 +175,4 @@ int DetectKnownPlatforms()
     {
         return PLATFORM_UNKNOWN;
     }
-}
+}

platforms/android/service/engine/jni/BinderComponent/TegraDetector.cpp

@@ -1,61 +0,0 @@
-#include "TegraDetector.h"
-#include <zlib.h>
-#include <string.h>
-
-#define KERNEL_CONFIG "/proc/config.gz"
-#define KERNEL_CONFIG_MAX_LINE_WIDTH 512
-#define KERNEL_CONFIG_TEGRA_MAGIC "CONFIG_ARCH_TEGRA=y"
-#define KERNEL_CONFIG_TEGRA2_MAGIC "CONFIG_ARCH_TEGRA_2x_SOC=y"
-#define KERNEL_CONFIG_TEGRA3_MAGIC "CONFIG_ARCH_TEGRA_3x_SOC=y"
-#define KERNEL_CONFIG_TEGRA4_MAGIC "CONFIG_ARCH_TEGRA_11x_SOC=y"
-#define MAX_DATA_LEN    4096
-
-int DetectTegra()
-{
-    int result = TEGRA_NOT_TEGRA;
-    gzFile kernelConfig = gzopen(KERNEL_CONFIG, "r");
-    if (kernelConfig != 0)
-    {
-        char tmpbuf[KERNEL_CONFIG_MAX_LINE_WIDTH];
-        const char *tegra_config = KERNEL_CONFIG_TEGRA_MAGIC;
-        const char *tegra2_config = KERNEL_CONFIG_TEGRA2_MAGIC;
-        const char *tegra3_config = KERNEL_CONFIG_TEGRA3_MAGIC;
-        const char *tegra4_config = KERNEL_CONFIG_TEGRA4_MAGIC;
-        int len = strlen(tegra_config);
-        int len2 = strlen(tegra2_config);
-        int len3 = strlen(tegra3_config);
-        int len4 = strlen(tegra4_config);
-        while (0 != gzgets(kernelConfig, tmpbuf, KERNEL_CONFIG_MAX_LINE_WIDTH))
-        {
-            if (0 == strncmp(tmpbuf, tegra_config, len))
-            {
-                result = 1;
-            }
-
-            if (0 == strncmp(tmpbuf, tegra2_config, len2))
-            {
-                result = 2;
-                break;
-            }
-
-            if (0 == strncmp(tmpbuf, tegra3_config, len3))
-            {
-                result = 3;
-                break;
-            }
-
-            if (0 == strncmp(tmpbuf, tegra4_config, len4))
-            {
-                result = 4;
-                break;
-            }
-        }
-        gzclose(kernelConfig);
-    }
-    else
-    {
-        result = TEGRA_DETECTOR_ERROR;
-    }
-
-    return result;
-}

platforms/android/service/engine/jni/BinderComponent/TegraDetector.h

@@ -2,8 +2,13 @@
 #define __TEGRA_DETECTOR_H__
 
 #define TEGRA_DETECTOR_ERROR -2
-#define TEGRA_NOT_TEGRA -1
+#define NOT_TEGRA -1
+#define TEGRA2  2
+#define TEGRA3  3
+#define TEGRA4i 4
+#define TEGRA4  5
+#define TEGRA5  6
 
 int DetectTegra();
 
-#endif
+#endif

platforms/android/service/engine/src/org/opencv/engine/HardwareDetector.java

@@ -30,11 +30,12 @@ public class HardwareDetector
     // GPU Acceleration options
     public static final int FEATURES_HAS_GPU = 0x010000;
 
-    public static final int PLATFORM_TEGRA  = 1;
-    public static final int PLATFORM_TEGRA2 = 2;
-    public static final int PLATFORM_TEGRA3 = 3;
-    public static final int PLATFORM_TEGRA4 = 4;
-
+    public static final int PLATFORM_TEGRA   = 1;
+    public static final int PLATFORM_TEGRA2  = 2;
+    public static final int PLATFORM_TEGRA3  = 3;
+    public static final int PLATFORM_TEGRA4i = 4;
+    public static final int PLATFORM_TEGRA4  = 5;
+    public static final int PLATFORM_TEGRA5  = 6;
 
     public static final int PLATFORM_UNKNOWN = 0;
 

platforms/android/service/engine/src/org/opencv/engine/manager/ManagerActivity.java

@@ -107,6 +107,10 @@ public class ManagerActivity extends Activity
             {
                 HardwarePlatformView.setText("Tegra 3");
             }
+            else if (HardwareDetector.PLATFORM_TEGRA4i == Platfrom)
+            {
+                HardwarePlatformView.setText("Tegra 4i");
+            }
             else
             {
                 HardwarePlatformView.setText("Tegra 4");

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']

platforms/scripts/camera_build.conf

@@ -21,3 +21,7 @@ native_camera_r4.2.0; armeabi-v7a; 14; /home/alexander/Projects/AndroidSource/4.
 native_camera_r4.2.0; armeabi;     14; /home/alexander/Projects/AndroidSource/4.2
 native_camera_r4.2.0; x86;         14; /home/alexander/Projects/AndroidSource/4.2
 native_camera_r4.2.0; mips;        14; /home/alexander/Projects/AndroidSource/4.2
+native_camera_r4.3.0; armeabi;     14; /home/alexander/Projects/AndroidSource/4.3
+native_camera_r4.3.0; armeabi-v7a; 14; /home/alexander/Projects/AndroidSource/4.3
+native_camera_r4.3.0; x86;         14; /home/alexander/Projects/AndroidSource/4.3
+native_camera_r4.3.0; mips;        14; /home/alexander/Projects/AndroidSource/4.3

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