From 31b7cdc695a060509399b092b837757e0a8f3f16 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sun, 2 Feb 2014 17:47:06 +0400 Subject: [PATCH] implemented OpenCL version of cv::updateMotionHistory --- modules/video/perf/opencl/perf_motempl.cpp | 36 +++++++++++++ modules/video/src/motempl.cpp | 44 +++++++++++++--- modules/video/src/opencl/updatemotionhistory.cl | 27 ++++++++++ modules/video/test/ocl/test_motempl.cpp | 67 +++++++++++++++++++++++++ 4 files changed, 166 insertions(+), 8 deletions(-) create mode 100644 modules/video/perf/opencl/perf_motempl.cpp create mode 100644 modules/video/src/opencl/updatemotionhistory.cl create mode 100644 modules/video/test/ocl/test_motempl.cpp diff --git a/modules/video/perf/opencl/perf_motempl.cpp b/modules/video/perf/opencl/perf_motempl.cpp new file mode 100644 index 0000000..7956857 --- /dev/null +++ b/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 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 diff --git a/modules/video/src/motempl.cpp b/modules/video/src/motempl.cpp index aa6d12d..3fc87e6 100644 --- a/modules/video/src/motempl.cpp +++ b/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 ) { - Mat silh = _silhouette.getMat(), mhi = _mhi.getMat(); + 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_Assert( silh.type() == CV_8U && mhi.type() == CV_32F ); - CV_Assert( silh.size() == mhi.size() ); + CV_OCL_RUN(_mhi.isUMat() && _mhi.dims() <= 2, + ocl_updateMotionHistory(_silhouette, _mhi, ts, delbound)) + Mat silh = _silhouette.getMat(), mhi = _mhi.getMat(); 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(y); float* mhiData = mhi.ptr(y); - x = 0; + int x = 0; #if CV_SSE2 if( useSIMD ) diff --git a/modules/video/src/opencl/updatemotionhistory.cl b/modules/video/src/opencl/updatemotionhistory.cl new file mode 100644 index 0000000..913e40b --- /dev/null +++ b/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; + } +} diff --git a/modules/video/test/ocl/test_motempl.cpp b/modules/video/test/ocl/test_motempl.cpp new file mode 100644 index 0000000..7b4c227 --- /dev/null +++ b/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 -- 2.7.4