From: Vladislav Vinogradov Date: Wed, 21 Sep 2011 07:02:45 +0000 (+0000) Subject: implemented gpu::copyMakeBorder for all border modes X-Git-Tag: accepted/2.0/20130307.220821~1837 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=64119dd9247f1aaeeb8afc4a974e5a4f0b178164;p=profile%2Fivi%2Fopencv.git implemented gpu::copyMakeBorder for all border modes --- diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index ffbc0da..dbb1107 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -662,8 +662,7 @@ namespace cv CV_EXPORTS void rotate(const GpuMat& src, GpuMat& dst, Size dsize, double angle, double xShift = 0, double yShift = 0, int interpolation = INTER_LINEAR, Stream& stream = Stream::Null()); //! copies 2D array to a larger destination array and pads borders with user-specifiable constant - //! supports CV_8UC1, CV_8UC4, CV_32SC1 and CV_32FC1 types - CV_EXPORTS void copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value = Scalar(), Stream& stream = Stream::Null()); + CV_EXPORTS void copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, const Scalar& value = Scalar(), Stream& stream = Stream::Null()); //! computes the integral image //! sum will have CV_32S type, but will contain unsigned int values diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp index b2acf4d..98e7edd 100644 --- a/modules/gpu/perf/perf_imgproc.cpp +++ b/modules/gpu/perf/perf_imgproc.cpp @@ -460,13 +460,15 @@ PERF_TEST_P(DevInfo_Size_MatType_Interpolation, rotate, testing::Combine(testing SANITY_CHECK(dst_host); } -PERF_TEST_P(DevInfo_Size_MatType, copyMakeBorder, testing::Combine(testing::ValuesIn(devices()), - testing::Values(GPU_TYPICAL_MAT_SIZES), - testing::Values(CV_8UC1, CV_8UC4, CV_32SC1))) +PERF_TEST_P(DevInfo_Size_MatType_BorderMode, copyMakeBorder, testing::Combine(testing::ValuesIn(devices()), + testing::Values(GPU_TYPICAL_MAT_SIZES), + testing::Values(CV_8UC1, CV_8UC4, CV_32FC1), + testing::Values((int)BORDER_REPLICATE, (int)BORDER_CONSTANT))) { DeviceInfo devInfo = std::tr1::get<0>(GetParam()); Size size = std::tr1::get<1>(GetParam()); int type = std::tr1::get<2>(GetParam()); + int borderType = std::tr1::get<3>(GetParam()); setDevice(devInfo.deviceID()); @@ -481,7 +483,7 @@ PERF_TEST_P(DevInfo_Size_MatType, copyMakeBorder, testing::Combine(testing::Valu SIMPLE_TEST_CYCLE() { - copyMakeBorder(src, dst, 5, 5, 5, 5); + copyMakeBorder(src, dst, 5, 5, 5, 5, borderType); } Mat dst_host(dst); diff --git a/modules/gpu/perf/perf_utility.hpp b/modules/gpu/perf/perf_utility.hpp index ec6b052..a57e367 100644 --- a/modules/gpu/perf/perf_utility.hpp +++ b/modules/gpu/perf/perf_utility.hpp @@ -49,6 +49,7 @@ typedef TestBaseWithParam< std::tr1::tuple > DevInfo typedef TestBaseWithParam< std::tr1::tuple > DevInfo_Size_MatType_NormType; typedef TestBaseWithParam< std::tr1::tuple > DevInfo_DescSize; typedef TestBaseWithParam< std::tr1::tuple > DevInfo_K_DescSize; +typedef TestBaseWithParam< std::tr1::tuple > DevInfo_Size_MatType_BorderMode; const cv::Size sz1800x1500 = cv::Size(1800, 1500); const cv::Size sz4700x3000 = cv::Size(4700, 3000); diff --git a/modules/gpu/src/cuda/column_filter.cu b/modules/gpu/src/cuda/column_filter.cu index 3c32769..b2eaa72 100644 --- a/modules/gpu/src/cuda/column_filter.cu +++ b/modules/gpu/src/cuda/column_filter.cu @@ -109,12 +109,6 @@ namespace cv { namespace gpu { namespace filters B b(src.rows); - if (!b.is_range_safe(-BLOCK_DIM_Y, (grid.y + 1) * BLOCK_DIM_Y - 1)) - { - cv::gpu::error("linearColumnFilter: can't use specified border extrapolation, image is too small, " - "try bigger image or another border extrapolation mode", __FILE__, __LINE__); - } - filter_krnls_column::linearColumnFilter<<>>(src, dst, anchor, b); cudaSafeCall( cudaGetLastError() ); diff --git a/modules/gpu/src/cuda/copy_make_border.cu b/modules/gpu/src/cuda/copy_make_border.cu new file mode 100644 index 0000000..8d70393 --- /dev/null +++ b/modules/gpu/src/cuda/copy_make_border.cu @@ -0,0 +1,127 @@ +/*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) 2000-2008, Intel Corporation, all rights reserved. +// Copyright (C) 2009, Willow Garage Inc., all rights reserved. +// Third party copyrights are property of their respective owners. +// +// 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*/ + +#include "internal_shared.hpp" +#include "opencv2/gpu/device/border_interpolate.hpp" + +using namespace cv::gpu; +using namespace cv::gpu::device; + +namespace cv { namespace gpu { namespace imgproc +{ + template __global__ void copyMakeBorder(const Ptr2D src, DevMem2D_ dst, int top, int left) + { + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (x < dst.cols && y < dst.rows) + dst.ptr(y)[x] = src(y - top, x - left); + } + + template