Merge pull request #2258 from ilya-lavrenov:tapi_video

modules/video/perf/opencl/perf_motempl.cpp

@@ -0,0 +1,36 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+// Copyright (C) 2014, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+
+#include "perf_precomp.hpp"
+#include "opencv2/ts/ocl_perf.hpp"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+///////////// UpdateMotionHistory ////////////////////////
+
+typedef TestBaseWithParam<Size> UpdateMotionHistoryFixture;
+
+OCL_PERF_TEST_P(UpdateMotionHistoryFixture, UpdateMotionHistory, OCL_TEST_SIZES)
+{
+    const Size size = GetParam();
+    checkDeviceMaxMemoryAllocSize(size, CV_32FC1);
+
+    UMat silhouette(size, CV_8UC1), mhi(size, CV_32FC1);
+    randu(silhouette, -5, 5);
+    declare.in(mhi, WARMUP_RNG);
+
+    OCL_TEST_CYCLE() cv::updateMotionHistory(silhouette, mhi, 1, 0.5);
+
+    SANITY_CHECK(mhi);
+}
+
+} } // namespace cvtest::ocl
+
+#endif // HAVE_OPENCL

modules/video/src/motempl.cpp

@@ -40,34 +40,62 @@
 //M*/
 
 #include "precomp.hpp"
+#include "opencl_kernels.hpp"
+
+#ifdef HAVE_OPENCL
+
+namespace  cv {
+
+static bool ocl_updateMotionHistory( InputArray _silhouette, InputOutputArray _mhi,
+                                     float timestamp, float delbound )
+{
+    ocl::Kernel k("updateMotionHistory", ocl::video::updatemotionhistory_oclsrc);
+    if (k.empty())
+        return false;
+
+    UMat silh = _silhouette.getUMat(), mhi = _mhi.getUMat();
+
+    k.args(ocl::KernelArg::ReadOnlyNoSize(silh), ocl::KernelArg::ReadWrite(mhi),
+           timestamp, delbound);
+
+    size_t globalsize[2] = { silh.cols, silh.rows };
+    return k.run(2, globalsize, NULL, false);
+}
+
+}
+
+#endif
 
 void cv::updateMotionHistory( InputArray _silhouette, InputOutputArray _mhi,
                               double timestamp, double duration )
 {
+    CV_Assert( _silhouette.type() == CV_8UC1 && _mhi.type() == CV_32FC1 );
+    CV_Assert( _silhouette.sameSize(_mhi) );
+
+    float ts = (float)timestamp;
+    float delbound = (float)(timestamp - duration);
+
+    CV_OCL_RUN(_mhi.isUMat() && _mhi.dims() <= 2,
+               ocl_updateMotionHistory(_silhouette, _mhi, ts, delbound))
+
     Mat silh = _silhouette.getMat(), mhi = _mhi.getMat();
-
-    CV_Assert( silh.type() == CV_8U && mhi.type() == CV_32F );
-    CV_Assert( silh.size() == mhi.size() );
-
     Size size = silh.size();
+
     if( silh.isContinuous() && mhi.isContinuous() )
     {
         size.width *= size.height;
         size.height = 1;
     }
 
-    float ts = (float)timestamp;
-    float delbound = (float)(timestamp - duration);
-    int x, y;
 #if CV_SSE2
     volatile bool useSIMD = cv::checkHardwareSupport(CV_CPU_SSE2);
 #endif
 
-    for( y = 0; y < size.height; y++ )
+    for(int y = 0; y < size.height; y++ )
     {
         const uchar* silhData = silh.ptr<uchar>(y);
         float* mhiData = mhi.ptr<float>(y);
-        x = 0;
+        int x = 0;
 
 #if CV_SSE2
         if( useSIMD )

modules/video/src/opencl/updatemotionhistory.cl

@@ -0,0 +1,27 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+// Copyright (C) 2014, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+
+__kernel void updateMotionHistory(__global const uchar * silh, int silh_step, int silh_offset,
+                                  __global uchar * mhiptr, int mhi_step, int mhi_offset, int mhi_rows, int mhi_cols,
+                                  float timestamp, float delbound)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < mhi_cols && y < mhi_rows)
+    {
+        int silh_index = mad24(y, silh_step, silh_offset + x);
+        int mhi_index = mad24(y, mhi_step, mhi_offset + x * (int)sizeof(float));
+
+        silh += silh_index;
+        __global float * mhi = (__global float *)(mhiptr + mhi_index);
+
+        float val = mhi[0];
+        val = silh[0] ? timestamp : val < delbound ? 0 : val;
+        mhi[0] = val;
+    }
+}

modules/video/test/ocl/test_motempl.cpp

@@ -0,0 +1,67 @@
+// This file is part of OpenCV project.
+// It is subject to the license terms in the LICENSE file found in the top-level directory
+// of this distribution and at http://opencv.org/license.html.
+
+// Copyright (C) 2014, Advanced Micro Devices, Inc., all rights reserved.
+// Third party copyrights are property of their respective owners.
+
+#include "test_precomp.hpp"
+#include "opencv2/ts/ocl_test.hpp"
+
+#ifdef HAVE_OPENCL
+
+namespace cvtest {
+namespace ocl {
+
+PARAM_TEST_CASE(UpdateMotionHistory, bool)
+{
+    double timestamp, duration;
+    bool use_roi;
+
+    TEST_DECLARE_INPUT_PARAMETER(silhouette)
+    TEST_DECLARE_OUTPUT_PARAMETER(mhi)
+
+    virtual void SetUp()
+    {
+        use_roi = GET_PARAM(0);
+    }
+
+    virtual void generateTestData()
+    {
+        Size roiSize = randomSize(1, MAX_VALUE);
+        Border silhouetteBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(silhouette, silhouette_roi, roiSize, silhouetteBorder, CV_8UC1, -11, 11);
+
+        Border mhiBorder = randomBorder(0, use_roi ? MAX_VALUE : 0);
+        randomSubMat(mhi, mhi_roi, roiSize, mhiBorder, CV_32FC1, 0, 1);
+
+        timestamp = randomDouble(0, 1);
+        duration = randomDouble(0, 1);
+        if (timestamp < duration)
+            std::swap(timestamp, duration);
+
+        UMAT_UPLOAD_INPUT_PARAMETER(silhouette)
+        UMAT_UPLOAD_OUTPUT_PARAMETER(mhi)
+    }
+};
+
+OCL_TEST_P(UpdateMotionHistory, Mat)
+{
+    for (int j = 0; j < test_loop_times; j++)
+    {
+        generateTestData();
+
+        OCL_OFF(cv::updateMotionHistory(silhouette_roi, mhi_roi, timestamp, duration));
+        OCL_ON(cv::updateMotionHistory(usilhouette_roi, umhi_roi, timestamp, duration));
+
+        OCL_EXPECT_MATS_NEAR(mhi, 0)
+    }
+}
+
+//////////////////////////////////////// Instantiation /////////////////////////////////////////
+
+OCL_INSTANTIATE_TEST_CASE_P(Video, UpdateMotionHistory, Values(false, true));
+
+} } // namespace cvtest::ocl
+
+#endif // HAVE_OPENCL