diff --git a/modules/gpu/include/opencv2/gpu.hpp b/modules/gpu/include/opencv2/gpu.hpp index 6605a58407..642c327a35 100644 --- a/modules/gpu/include/opencv2/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu.hpp @@ -168,10 +168,6 @@ CV_EXPORTS void buildWarpSphericalMaps(Size src_size, Rect dst_roi, const Mat &K 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 -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 //! supports only CV_8UC1 source type diff --git a/modules/gpu/perf/perf_imgproc.cpp b/modules/gpu/perf/perf_imgproc.cpp index 9a1168a521..02a1e6cd46 100644 --- a/modules/gpu/perf/perf_imgproc.cpp +++ b/modules/gpu/perf/perf_imgproc.cpp @@ -325,46 +325,6 @@ PERF_TEST_P(Sz_Depth_Cn_Inter_Border, ImgProc_WarpPerspective, } } -////////////////////////////////////////////////////////////////////// -// CopyMakeBorder - -DEF_PARAM_TEST(Sz_Depth_Cn_Border, cv::Size, MatDepth, MatCn, BorderMode); - -PERF_TEST_P(Sz_Depth_Cn_Border, ImgProc_CopyMakeBorder, - Combine(GPU_TYPICAL_MAT_SIZES, - Values(CV_8U, CV_16U, CV_32F), - GPU_CHANNELS_1_3_4, - ALL_BORDER_MODES)) -{ - const cv::Size size = GET_PARAM(0); - const int depth = GET_PARAM(1); - const int channels = GET_PARAM(2); - const int borderMode = GET_PARAM(3); - - const int type = CV_MAKE_TYPE(depth, channels); - - cv::Mat src(size, type); - declare.in(src, WARMUP_RNG); - - if (PERF_RUN_GPU()) - { - const cv::gpu::GpuMat d_src(src); - cv::gpu::GpuMat dst; - - TEST_CYCLE() cv::gpu::copyMakeBorder(d_src, dst, 5, 5, 5, 5, borderMode); - - GPU_SANITY_CHECK(dst); - } - else - { - cv::Mat dst; - - TEST_CYCLE() cv::copyMakeBorder(src, dst, 5, 5, 5, 5, borderMode); - - CPU_SANITY_CHECK(dst); - } -} - ////////////////////////////////////////////////////////////////////// // Threshold diff --git a/modules/gpu/src/imgproc.cpp b/modules/gpu/src/imgproc.cpp index a3c9694fa6..54a6937006 100644 --- a/modules/gpu/src/imgproc.cpp +++ b/modules/gpu/src/imgproc.cpp @@ -51,7 +51,6 @@ void cv::gpu::meanShiftFiltering(const GpuMat&, GpuMat&, int, int, TermCriteria, void cv::gpu::meanShiftProc(const GpuMat&, GpuMat&, GpuMat&, int, int, TermCriteria, Stream&) { throw_no_cuda(); } void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, Stream&) { throw_no_cuda(); } void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, int, Stream&) { throw_no_cuda(); } -void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, int, const Scalar&, Stream&) { throw_no_cuda(); } void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_no_cuda(); } @@ -235,115 +234,6 @@ void cv::gpu::reprojectImageTo3D(const GpuMat& disp, GpuMat& xyz, const Mat& Q, funcs[dst_cn == 4][disp.type()](disp, xyz, Q.ptr(), StreamAccessor::getStream(stream)); } -//////////////////////////////////////////////////////////////////////// -// copyMakeBorder - -namespace cv { namespace gpu { namespace cudev -{ - namespace imgproc - { - template void copyMakeBorder_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderMode, const T* borderValue, cudaStream_t stream); - } -}}} - -namespace -{ - template void copyMakeBorder_caller(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream) - { - using namespace ::cv::gpu::cudev::imgproc; - - Scalar_ val(saturate_cast(value[0]), saturate_cast(value[1]), saturate_cast(value[2]), saturate_cast(value[3])); - - copyMakeBorder_gpu(src, dst, top, left, borderType, val.val, stream); - } -} - -#if defined __GNUC__ && __GNUC__ > 2 && __GNUC_MINOR__ > 4 -typedef Npp32s __attribute__((__may_alias__)) Npp32s_a; -#else -typedef Npp32s Npp32s_a; -#endif - -void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, const Scalar& value, Stream& s) -{ - CV_Assert(src.depth() <= CV_32F && src.channels() <= 4); - CV_Assert(borderType == BORDER_REFLECT101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP); - - dst.create(src.rows + top + bottom, src.cols + left + right, src.type()); - - cudaStream_t stream = StreamAccessor::getStream(s); - - if (borderType == BORDER_CONSTANT && (src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1)) - { - NppiSize srcsz; - srcsz.width = src.cols; - srcsz.height = src.rows; - - NppiSize dstsz; - dstsz.width = dst.cols; - dstsz.height = dst.rows; - - NppStreamHandler h(stream); - - switch (src.type()) - { - case CV_8UC1: - { - Npp8u nVal = saturate_cast(value[0]); - nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr(), static_cast(src.step), srcsz, - dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); - break; - } - case CV_8UC4: - { - Npp8u nVal[] = {saturate_cast(value[0]), saturate_cast(value[1]), saturate_cast(value[2]), saturate_cast(value[3])}; - nppSafeCall( nppiCopyConstBorder_8u_C4R(src.ptr(), static_cast(src.step), srcsz, - dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); - break; - } - case CV_32SC1: - { - Npp32s nVal = saturate_cast(value[0]); - nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr(), static_cast(src.step), srcsz, - dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); - break; - } - case CV_32FC1: - { - Npp32f val = saturate_cast(value[0]); - Npp32s nVal = *(reinterpret_cast(&val)); - nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr(), static_cast(src.step), srcsz, - dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); - break; - } - } - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } - else - { - typedef void (*caller_t)(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream); - static const caller_t callers[6][4] = - { - { copyMakeBorder_caller , copyMakeBorder_caller , copyMakeBorder_caller , copyMakeBorder_caller}, - {0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/ , 0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/}, - { copyMakeBorder_caller , 0/*copyMakeBorder_caller*/, copyMakeBorder_caller , copyMakeBorder_caller}, - { copyMakeBorder_caller , 0/*copyMakeBorder_caller*/ , copyMakeBorder_caller , copyMakeBorder_caller}, - {0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/ , 0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/}, - { copyMakeBorder_caller , 0/*copyMakeBorder_caller*/ , copyMakeBorder_caller , copyMakeBorder_caller} - }; - - caller_t func = callers[src.depth()][src.channels() - 1]; - CV_Assert(func != 0); - - int gpuBorderType; - CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType)); - - func(src, dst, top, left, gpuBorderType, value, stream); - } -} - ////////////////////////////////////////////////////////////////////////////// // buildWarpPlaneMaps diff --git a/modules/gpu/test/test_copy_make_border.cpp b/modules/gpu/test/test_copy_make_border.cpp deleted file mode 100644 index 24a75c0235..0000000000 --- a/modules/gpu/test/test_copy_make_border.cpp +++ /dev/null @@ -1,106 +0,0 @@ -/*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 "test_precomp.hpp" - -#ifdef HAVE_CUDA - -using namespace cvtest; - -namespace -{ - IMPLEMENT_PARAM_CLASS(Border, int) -} - -PARAM_TEST_CASE(CopyMakeBorder, cv::gpu::DeviceInfo, cv::Size, MatType, Border, BorderType, UseRoi) -{ - cv::gpu::DeviceInfo devInfo; - cv::Size size; - int type; - int border; - int borderType; - bool useRoi; - - virtual void SetUp() - { - devInfo = GET_PARAM(0); - size = GET_PARAM(1); - type = GET_PARAM(2); - border = GET_PARAM(3); - borderType = GET_PARAM(4); - useRoi = GET_PARAM(5); - - cv::gpu::setDevice(devInfo.deviceID()); - } -}; - -GPU_TEST_P(CopyMakeBorder, Accuracy) -{ - cv::Mat src = randomMat(size, type); - cv::Scalar val = randomScalar(0, 255); - - cv::gpu::GpuMat dst = createMat(cv::Size(size.width + 2 * border, size.height + 2 * border), type, useRoi); - cv::gpu::copyMakeBorder(loadMat(src, useRoi), dst, border, border, border, border, borderType, val); - - cv::Mat dst_gold; - cv::copyMakeBorder(src, dst_gold, border, border, border, border, borderType, val); - - EXPECT_MAT_NEAR(dst_gold, dst, 0.0); -} - -INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CopyMakeBorder, testing::Combine( - ALL_DEVICES, - DIFFERENT_SIZES, - testing::Values(MatType(CV_8UC1), - MatType(CV_8UC3), - MatType(CV_8UC4), - MatType(CV_16UC1), - MatType(CV_16UC3), - MatType(CV_16UC4), - MatType(CV_32FC1), - MatType(CV_32FC3), - MatType(CV_32FC4)), - testing::Values(Border(1), Border(10), Border(50)), - ALL_BORDER_TYPES, - WHOLE_SUBMAT)); - -#endif // HAVE_CUDA diff --git a/modules/gpuarithm/CMakeLists.txt b/modules/gpuarithm/CMakeLists.txt index 99a6fcce13..04a6b2cc20 100644 --- a/modules/gpuarithm/CMakeLists.txt +++ b/modules/gpuarithm/CMakeLists.txt @@ -6,7 +6,7 @@ set(the_description "GPU-accelerated Operations on Matrices") ocv_warnings_disable(CMAKE_CXX_FLAGS -Wundef -Wmissing-declarations) -ocv_define_module(gpuarithm opencv_core) +ocv_define_module(gpuarithm opencv_core OPTIONAL opencv_imgproc) if(HAVE_CUBLAS) CUDA_ADD_CUBLAS_TO_TARGET(${the_module}) diff --git a/modules/gpuarithm/include/opencv2/gpuarithm.hpp b/modules/gpuarithm/include/opencv2/gpuarithm.hpp index 5724cd53c7..03458ea045 100644 --- a/modules/gpuarithm/include/opencv2/gpuarithm.hpp +++ b/modules/gpuarithm/include/opencv2/gpuarithm.hpp @@ -279,6 +279,10 @@ CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh, doubl //! output will have CV_32FC1 type CV_EXPORTS void rectStdDev(const GpuMat& src, const GpuMat& sqr, GpuMat& dst, const Rect& rect, Stream& stream = Stream::Null()); +//! copies 2D array to a larger destination array and pads borders with user-specifiable constant +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()); + }} // namespace cv { namespace gpu { #endif /* __OPENCV_GPUARITHM_HPP__ */ diff --git a/modules/gpuarithm/perf/perf_core.cpp b/modules/gpuarithm/perf/perf_core.cpp index 829637f4e3..a9a6e360b5 100644 --- a/modules/gpuarithm/perf/perf_core.cpp +++ b/modules/gpuarithm/perf/perf_core.cpp @@ -2155,3 +2155,47 @@ PERF_TEST_P(Sz_Depth_NormType, Core_Normalize, CPU_SANITY_CHECK(dst); } } + +////////////////////////////////////////////////////////////////////// +// CopyMakeBorder + +#ifdef HAVE_OPENCV_IMGPROC + +DEF_PARAM_TEST(Sz_Depth_Cn_Border, cv::Size, MatDepth, MatCn, BorderMode); + +PERF_TEST_P(Sz_Depth_Cn_Border, ImgProc_CopyMakeBorder, + Combine(GPU_TYPICAL_MAT_SIZES, + Values(CV_8U, CV_16U, CV_32F), + GPU_CHANNELS_1_3_4, + ALL_BORDER_MODES)) +{ + const cv::Size size = GET_PARAM(0); + const int depth = GET_PARAM(1); + const int channels = GET_PARAM(2); + const int borderMode = GET_PARAM(3); + + const int type = CV_MAKE_TYPE(depth, channels); + + cv::Mat src(size, type); + declare.in(src, WARMUP_RNG); + + if (PERF_RUN_GPU()) + { + const cv::gpu::GpuMat d_src(src); + cv::gpu::GpuMat dst; + + TEST_CYCLE() cv::gpu::copyMakeBorder(d_src, dst, 5, 5, 5, 5, borderMode); + + GPU_SANITY_CHECK(dst); + } + else + { + cv::Mat dst; + + TEST_CYCLE() cv::copyMakeBorder(src, dst, 5, 5, 5, 5, borderMode); + + CPU_SANITY_CHECK(dst); + } +} + +#endif diff --git a/modules/gpuarithm/perf/perf_precomp.hpp b/modules/gpuarithm/perf/perf_precomp.hpp index 06bc20b9b9..bee378064b 100644 --- a/modules/gpuarithm/perf/perf_precomp.hpp +++ b/modules/gpuarithm/perf/perf_precomp.hpp @@ -57,6 +57,12 @@ #include "opencv2/core.hpp" #include "opencv2/gpuarithm.hpp" +#include "opencv2/opencv_modules.hpp" + +#ifdef HAVE_OPENCV_IMGPROC +# include "opencv2/imgproc.hpp" +#endif + #ifdef GTEST_CREATE_SHARED_LIBRARY #error no modules except ts should have GTEST_CREATE_SHARED_LIBRARY defined #endif diff --git a/modules/gpuarithm/src/arithm.cpp b/modules/gpuarithm/src/arithm.cpp index d452e3ae74..40242876d4 100644 --- a/modules/gpuarithm/src/arithm.cpp +++ b/modules/gpuarithm/src/arithm.cpp @@ -60,6 +60,7 @@ void cv::gpu::cartToPolar(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, void cv::gpu::polarToCart(const GpuMat&, const GpuMat&, GpuMat&, GpuMat&, bool, Stream&) { throw_no_cuda(); } void cv::gpu::normalize(const GpuMat&, GpuMat&, double, double, int, int, const GpuMat&) { throw_no_cuda(); } void cv::gpu::normalize(const GpuMat&, GpuMat&, double, double, int, int, const GpuMat&, GpuMat&, GpuMat&) { throw_no_cuda(); } +void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, int, const Scalar&, Stream&) { throw_no_cuda(); } #else /* !defined (HAVE_CUDA) */ @@ -608,4 +609,113 @@ void cv::gpu::normalize(const GpuMat& src, GpuMat& dst, double a, double b, int } } +//////////////////////////////////////////////////////////////////////// +// copyMakeBorder + +namespace cv { namespace gpu { namespace cudev +{ + namespace imgproc + { + template void copyMakeBorder_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderMode, const T* borderValue, cudaStream_t stream); + } +}}} + +namespace +{ + template void copyMakeBorder_caller(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream) + { + using namespace ::cv::gpu::cudev::imgproc; + + Scalar_ val(saturate_cast(value[0]), saturate_cast(value[1]), saturate_cast(value[2]), saturate_cast(value[3])); + + copyMakeBorder_gpu(src, dst, top, left, borderType, val.val, stream); + } +} + +#if defined __GNUC__ && __GNUC__ > 2 && __GNUC_MINOR__ > 4 +typedef Npp32s __attribute__((__may_alias__)) Npp32s_a; +#else +typedef Npp32s Npp32s_a; +#endif + +void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, const Scalar& value, Stream& s) +{ + CV_Assert(src.depth() <= CV_32F && src.channels() <= 4); + CV_Assert(borderType == IPL_BORDER_REFLECT_101 || borderType == IPL_BORDER_REPLICATE || borderType == IPL_BORDER_CONSTANT || borderType == IPL_BORDER_REFLECT || borderType == IPL_BORDER_WRAP); + + dst.create(src.rows + top + bottom, src.cols + left + right, src.type()); + + cudaStream_t stream = StreamAccessor::getStream(s); + + if (borderType == IPL_BORDER_CONSTANT && (src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1)) + { + NppiSize srcsz; + srcsz.width = src.cols; + srcsz.height = src.rows; + + NppiSize dstsz; + dstsz.width = dst.cols; + dstsz.height = dst.rows; + + NppStreamHandler h(stream); + + switch (src.type()) + { + case CV_8UC1: + { + Npp8u nVal = saturate_cast(value[0]); + nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr(), static_cast(src.step), srcsz, + dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); + break; + } + case CV_8UC4: + { + Npp8u nVal[] = {saturate_cast(value[0]), saturate_cast(value[1]), saturate_cast(value[2]), saturate_cast(value[3])}; + nppSafeCall( nppiCopyConstBorder_8u_C4R(src.ptr(), static_cast(src.step), srcsz, + dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); + break; + } + case CV_32SC1: + { + Npp32s nVal = saturate_cast(value[0]); + nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr(), static_cast(src.step), srcsz, + dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); + break; + } + case CV_32FC1: + { + Npp32f val = saturate_cast(value[0]); + Npp32s nVal = *(reinterpret_cast(&val)); + nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr(), static_cast(src.step), srcsz, + dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); + break; + } + } + + if (stream == 0) + cudaSafeCall( cudaDeviceSynchronize() ); + } + else + { + typedef void (*caller_t)(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream); + static const caller_t callers[6][4] = + { + { copyMakeBorder_caller , copyMakeBorder_caller , copyMakeBorder_caller , copyMakeBorder_caller}, + {0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/ , 0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/}, + { copyMakeBorder_caller , 0/*copyMakeBorder_caller*/, copyMakeBorder_caller , copyMakeBorder_caller}, + { copyMakeBorder_caller , 0/*copyMakeBorder_caller*/ , copyMakeBorder_caller , copyMakeBorder_caller}, + {0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/ , 0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/}, + { copyMakeBorder_caller , 0/*copyMakeBorder_caller*/ , copyMakeBorder_caller , copyMakeBorder_caller} + }; + + caller_t func = callers[src.depth()][src.channels() - 1]; + CV_Assert(func != 0); + + int gpuBorderType; + CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType)); + + func(src, dst, top, left, gpuBorderType, value, stream); + } +} + #endif /* !defined (HAVE_CUDA) */ diff --git a/modules/gpu/src/cuda/copy_make_border.cu b/modules/gpuarithm/src/cuda/copy_make_border.cu similarity index 100% rename from modules/gpu/src/cuda/copy_make_border.cu rename to modules/gpuarithm/src/cuda/copy_make_border.cu diff --git a/modules/gpuarithm/test/test_core.cpp b/modules/gpuarithm/test/test_core.cpp index b2072c25fc..613d7122ac 100644 --- a/modules/gpuarithm/test/test_core.cpp +++ b/modules/gpuarithm/test/test_core.cpp @@ -3607,4 +3607,68 @@ INSTANTIATE_TEST_CASE_P(GPU_Core, Normalize, testing::Combine( testing::Values(NormCode(cv::NORM_L1), NormCode(cv::NORM_L2), NormCode(cv::NORM_INF), NormCode(cv::NORM_MINMAX)), WHOLE_SUBMAT)); +////////////////////////////////////////////////////////////////////////////// +// CopyMakeBorder + +#ifdef HAVE_OPENCV_IMGPROC + +namespace +{ + IMPLEMENT_PARAM_CLASS(Border, int) +} + +PARAM_TEST_CASE(CopyMakeBorder, cv::gpu::DeviceInfo, cv::Size, MatType, Border, BorderType, UseRoi) +{ + cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; + int border; + int borderType; + bool useRoi; + + virtual void SetUp() + { + devInfo = GET_PARAM(0); + size = GET_PARAM(1); + type = GET_PARAM(2); + border = GET_PARAM(3); + borderType = GET_PARAM(4); + useRoi = GET_PARAM(5); + + cv::gpu::setDevice(devInfo.deviceID()); + } +}; + +GPU_TEST_P(CopyMakeBorder, Accuracy) +{ + cv::Mat src = randomMat(size, type); + cv::Scalar val = randomScalar(0, 255); + + cv::gpu::GpuMat dst = createMat(cv::Size(size.width + 2 * border, size.height + 2 * border), type, useRoi); + cv::gpu::copyMakeBorder(loadMat(src, useRoi), dst, border, border, border, border, borderType, val); + + cv::Mat dst_gold; + cv::copyMakeBorder(src, dst_gold, border, border, border, border, borderType, val); + + EXPECT_MAT_NEAR(dst_gold, dst, 0.0); +} + +INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CopyMakeBorder, testing::Combine( + ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), + MatType(CV_8UC3), + MatType(CV_8UC4), + MatType(CV_16UC1), + MatType(CV_16UC3), + MatType(CV_16UC4), + MatType(CV_32FC1), + MatType(CV_32FC3), + MatType(CV_32FC4)), + testing::Values(Border(1), Border(10), Border(50)), + ALL_BORDER_TYPES, + WHOLE_SUBMAT)); + +#endif + #endif // HAVE_CUDA diff --git a/modules/gpuarithm/test/test_precomp.hpp b/modules/gpuarithm/test/test_precomp.hpp index 0896277583..800ed31c07 100644 --- a/modules/gpuarithm/test/test_precomp.hpp +++ b/modules/gpuarithm/test/test_precomp.hpp @@ -57,4 +57,10 @@ #include "opencv2/core.hpp" #include "opencv2/gpuarithm.hpp" +#include "opencv2/opencv_modules.hpp" + +#ifdef HAVE_OPENCV_IMGPROC +# include "opencv2/imgproc.hpp" +#endif + #endif