From d2ffd8ed211c3ce814e7b1f37faee67080844596 Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Mon, 20 Jan 2014 18:35:32 +0400 Subject: [PATCH] implemented OpenCL version of cv::preCornerDetect --- modules/imgproc/perf/opencl/perf_imgproc.cpp | 19 +++++ modules/imgproc/src/corner.cpp | 43 ++++++++++- modules/imgproc/src/opencl/precornerdetect.cl | 75 +++++++++++++++++++ modules/imgproc/test/ocl/test_imgproc.cpp | 28 ++++++- 4 files changed, 160 insertions(+), 5 deletions(-) create mode 100644 modules/imgproc/src/opencl/precornerdetect.cl diff --git a/modules/imgproc/perf/opencl/perf_imgproc.cpp b/modules/imgproc/perf/opencl/perf_imgproc.cpp index 6f0e463654..8d4b5c2cd3 100644 --- a/modules/imgproc/perf/opencl/perf_imgproc.cpp +++ b/modules/imgproc/perf/opencl/perf_imgproc.cpp @@ -132,6 +132,25 @@ OCL_PERF_TEST_P(CornerHarrisFixture, CornerHarris, SANITY_CHECK(dst, 5e-6, ERROR_RELATIVE); } +///////////// PreCornerDetect //////////////////////// + +typedef Size_MatType PreCornerDetectFixture; + +OCL_PERF_TEST_P(PreCornerDetectFixture, PreCornerDetect, + ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_8UC1, CV_32FC1))) +{ + const Size_MatType_t params = GetParam(); + const Size srcSize = get<0>(params); + const int type = get<1>(params), borderType = BORDER_REFLECT; + + UMat src(srcSize, type), dst(srcSize, CV_32FC1); + declare.in(src, WARMUP_RNG).out(dst); + + OCL_TEST_CYCLE() cv::preCornerDetect(src, dst, 3, borderType); + + SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE); +} + ///////////// Integral //////////////////////// typedef tuple IntegralParams; diff --git a/modules/imgproc/src/corner.cpp b/modules/imgproc/src/corner.cpp index e5d54c682c..fe76481911 100644 --- a/modules/imgproc/src/corner.cpp +++ b/modules/imgproc/src/corner.cpp @@ -363,6 +363,36 @@ static bool ocl_cornerMinEigenValVecs(InputArray _src, OutputArray _dst, int blo return cornelKernel.run(2, globalsize, localsize, false); } +static bool ocl_preCornerDetect( InputArray _src, OutputArray _dst, int ksize, int borderType, int depth ) +{ + UMat Dx, Dy, D2x, D2y, Dxy; + + Sobel( _src, Dx, CV_32F, 1, 0, ksize, 1, 0, borderType ); + Sobel( _src, Dy, CV_32F, 0, 1, ksize, 1, 0, borderType ); + Sobel( _src, D2x, CV_32F, 2, 0, ksize, 1, 0, borderType ); + Sobel( _src, D2y, CV_32F, 0, 2, ksize, 1, 0, borderType ); + Sobel( _src, Dxy, CV_32F, 1, 1, ksize, 1, 0, borderType ); + + _dst.create( _src.size(), CV_32FC1 ); + UMat dst = _dst.getUMat(); + + double factor = 1 << (ksize - 1); + if( depth == CV_8U ) + factor *= 255; + factor = 1./(factor * factor * factor); + + ocl::Kernel k("preCornerDetect", ocl::imgproc::precornerdetect_oclsrc); + if (k.empty()) + return false; + + k.args(ocl::KernelArg::ReadOnlyNoSize(Dx), ocl::KernelArg::ReadOnlyNoSize(Dy), + ocl::KernelArg::ReadOnlyNoSize(D2x), ocl::KernelArg::ReadOnlyNoSize(D2y), + ocl::KernelArg::ReadOnlyNoSize(Dxy), ocl::KernelArg::WriteOnly(dst), (float)factor); + + size_t globalsize[2] = { dst.cols, dst.rows }; + return k.run(2, globalsize, NULL, false); +} + } void cv::cornerMinEigenVal( InputArray _src, OutputArray _dst, int blockSize, int ksize, int borderType ) @@ -381,7 +411,7 @@ void cv::cornerMinEigenVal( InputArray _src, OutputArray _dst, int blockSize, in void cv::cornerHarris( InputArray _src, OutputArray _dst, int blockSize, int ksize, double k, int borderType ) { if (ocl::useOpenCL() && _src.dims() <= 2 && _dst.isUMat() && - ocl_cornerMinEigenValVecs(_src, _dst, blockSize, ksize, k, borderType, HARRIS)) + ocl_cornerMinEigenValVecs(_src, _dst, blockSize, ksize, k, borderType, HARRIS)) return; Mat src = _src.getMat(); @@ -406,10 +436,15 @@ void cv::cornerEigenValsAndVecs( InputArray _src, OutputArray _dst, int blockSiz void cv::preCornerDetect( InputArray _src, OutputArray _dst, int ksize, int borderType ) { - Mat Dx, Dy, D2x, D2y, Dxy, src = _src.getMat(); + int type = _src.type(); + CV_Assert( type == CV_8UC1 || type == CV_32FC1 ); - CV_Assert( src.type() == CV_8UC1 || src.type() == CV_32FC1 ); - _dst.create( src.size(), CV_32F ); + if (ocl::useOpenCL() && _src.dims() <= 2 && _dst.isUMat() && + ocl_preCornerDetect(_src, _dst, ksize, borderType, CV_MAT_DEPTH(type))) + return; + + Mat Dx, Dy, D2x, D2y, Dxy, src = _src.getMat(); + _dst.create( src.size(), CV_32FC1 ); Mat dst = _dst.getMat(); Sobel( src, Dx, CV_32F, 1, 0, ksize, 1, 0, borderType ); diff --git a/modules/imgproc/src/opencl/precornerdetect.cl b/modules/imgproc/src/opencl/precornerdetect.cl new file mode 100644 index 0000000000..32e6c2ce53 --- /dev/null +++ b/modules/imgproc/src/opencl/precornerdetect.cl @@ -0,0 +1,75 @@ +/*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, Institute Of Software Chinese Academy Of Science, all rights reserved. +// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// @Authors +// Shengen Yan,yanshengen@gmail.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 materials 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*/ + +__kernel void preCornerDetect(__global const uchar * Dxptr, int dx_step, int dx_offset, + __global const uchar * Dyptr, int dy_step, int dy_offset, + __global const uchar * D2xptr, int d2x_step, int d2x_offset, + __global const uchar * D2yptr, int d2y_step, int d2y_offset, + __global const uchar * Dxyptr, int dxy_step, int dxy_offset, + __global uchar * dstptr, int dst_step, int dst_offset, + int dst_rows, int dst_cols, float factor) +{ + int x = get_global_id(0); + int y = get_global_id(1); + + if (x < dst_cols && y < dst_rows) + { + int dx_index = mad24(dx_step, y, (int)sizeof(float) * x + dx_offset); + int dy_index = mad24(dy_step, y, (int)sizeof(float) * x + dy_offset); + int d2x_index = mad24(d2x_step, y, (int)sizeof(float) * x + d2x_offset); + int d2y_index = mad24(d2y_step, y, (int)sizeof(float) * x + d2y_offset); + int dxy_index = mad24(dxy_step, y, (int)sizeof(float) * x + dxy_offset); + int dst_index = mad24(dst_step, y, (int)sizeof(float) * x + dst_offset); + + float dx = *(__global const float *)(Dxptr + dx_index); + float dy = *(__global const float *)(Dyptr + dy_index); + float d2x = *(__global const float *)(D2xptr + d2x_index); + float d2y = *(__global const float *)(D2yptr + d2y_index); + float dxy = *(__global const float *)(Dxyptr + dxy_index); + __global float * dst = (__global float *)(dstptr + dst_index); + + dst[0] = factor * (dx*dx*d2y + dy*dy*d2x - 2*dx*dy*dxy); + } +} diff --git a/modules/imgproc/test/ocl/test_imgproc.cpp b/modules/imgproc/test/ocl/test_imgproc.cpp index 2d34d71549..78b2e573d4 100644 --- a/modules/imgproc/test/ocl/test_imgproc.cpp +++ b/modules/imgproc/test/ocl/test_imgproc.cpp @@ -259,6 +259,26 @@ OCL_TEST_P(CornerHarris, Mat) } } +//////////////////////////////// preCornerDetect ////////////////////////////////////////// + +typedef ImgprocTestBase PreCornerDetect; + +OCL_TEST_P(PreCornerDetect, Mat) +{ + for (int j = 0; j < test_loop_times; j++) + { + random_roi(); + + const int apertureSize = blockSize; + + OCL_OFF(cv::preCornerDetect(src_roi, dst_roi, apertureSize, borderType)); + OCL_ON(cv::preCornerDetect(usrc_roi, udst_roi, apertureSize, borderType)); + + Near(1e-6, true); + } +} + + ////////////////////////////////// integral ///////////////////////////////////////////////// struct Integral : @@ -341,7 +361,6 @@ struct Threshold : virtual void SetUp() { type = GET_PARAM(0); - blockSize = GET_PARAM(1); thresholdType = GET_PARAM(2); useRoi = GET_PARAM(3); } @@ -437,6 +456,13 @@ OCL_INSTANTIATE_TEST_CASE_P(Imgproc, CornerHarris, Combine( (BorderType)BORDER_REFLECT, (BorderType)BORDER_REFLECT_101), Bool())); +OCL_INSTANTIATE_TEST_CASE_P(Imgproc, PreCornerDetect, Combine( + Values((MatType)CV_8UC1, CV_32FC1), + Values(3, 5), + Values( (BorderType)BORDER_CONSTANT, (BorderType)BORDER_REPLICATE, + (BorderType)BORDER_REFLECT, (BorderType)BORDER_REFLECT_101), + Bool())); + OCL_INSTANTIATE_TEST_CASE_P(Imgproc, Integral, Combine( Values((MatType)CV_8UC1), // TODO does not work with CV_32F, CV_64F Values(CV_32SC1, CV_32FC1), // desired sdepth