mirror of https://github.com/opencv/opencv.git
Merge branch 'master' of https://github.com/Itseez/opencv into brief_cl
Matthias Bady
11 years ago
commit
67e2d7ec43
82 changed files with 2508 additions and 436 deletions
@ -1,3 +1,3 @@ |
||||
#include "test_precomp.hpp" |
||||
|
||||
CV_TEST_MAIN("cv") |
||||
CV_TEST_MAIN("") |
||||
|
||||
@ -0,0 +1,82 @@ |
||||
/*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, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Fangfang Bai, fangfang@multicorewareinc.com
|
||||
// Jin Ma, jin@multicorewareinc.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*/
|
||||
|
||||
#include "perf_precomp.hpp" |
||||
#include "opencv2/ts/ocl_perf.hpp" |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
namespace cvtest { |
||||
namespace ocl { |
||||
|
||||
///////////// BlendLinear ////////////////////////
|
||||
|
||||
typedef Size_MatType BlendLinearFixture; |
||||
|
||||
OCL_PERF_TEST_P(BlendLinearFixture, BlendLinear, ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32FC1, CV_32FC4))) |
||||
{ |
||||
Size_MatType_t params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int srcType = get<1>(params); |
||||
const double eps = CV_MAT_DEPTH(srcType) <= CV_32S ? 1.0 : 0.2; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, srcType); |
||||
|
||||
UMat src1(srcSize, srcType), src2(srcSize, srcType), dst(srcSize, srcType); |
||||
UMat weights1(srcSize, CV_32FC1), weights2(srcSize, CV_32FC1); |
||||
|
||||
declare.in(src1, src2, WARMUP_RNG).in(weights1, weights2, WARMUP_READ).out(dst); |
||||
randu(weights1, 0, 1); |
||||
randu(weights2, 0, 1); |
||||
|
||||
OCL_TEST_CYCLE() cv::blendLinear(src1, src2, weights1, weights2, dst); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
||||
@ -0,0 +1,104 @@ |
||||
/*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, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Fangfang Bai, fangfang@multicorewareinc.com
|
||||
// Jin Ma, jin@multicorewareinc.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*/
|
||||
|
||||
#include "perf_precomp.hpp" |
||||
#include "opencv2/ts/ocl_perf.hpp" |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
namespace cvtest { |
||||
namespace ocl { |
||||
|
||||
using std::tr1::make_tuple; |
||||
|
||||
///////////// cvtColor////////////////////////
|
||||
|
||||
CV_ENUM(ConversionTypes, COLOR_RGB2GRAY, COLOR_RGB2BGR, COLOR_RGB2YUV, COLOR_YUV2RGB, COLOR_RGB2YCrCb, |
||||
COLOR_YCrCb2RGB, COLOR_RGB2XYZ, COLOR_XYZ2RGB, COLOR_RGB2HSV, COLOR_HSV2RGB, COLOR_RGB2HLS, |
||||
COLOR_HLS2RGB, COLOR_BGR5652BGR, COLOR_BGR2BGR565, COLOR_RGBA2mRGBA, COLOR_mRGBA2RGBA, COLOR_YUV2RGB_NV12) |
||||
|
||||
typedef tuple<Size, tuple<ConversionTypes, int, int> > CvtColorParams; |
||||
typedef TestBaseWithParam<CvtColorParams> CvtColorFixture; |
||||
|
||||
OCL_PERF_TEST_P(CvtColorFixture, CvtColor, testing::Combine( |
||||
OCL_TEST_SIZES, |
||||
testing::Values( |
||||
make_tuple(ConversionTypes(COLOR_RGB2GRAY), 3, 1), |
||||
make_tuple(ConversionTypes(COLOR_RGB2BGR), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_RGB2YUV), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_YUV2RGB), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_RGB2YCrCb), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_YCrCb2RGB), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_RGB2XYZ), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_XYZ2RGB), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_RGB2HSV), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_HSV2RGB), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_RGB2HLS), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_HLS2RGB), 3, 3), |
||||
make_tuple(ConversionTypes(COLOR_BGR5652BGR), 2, 3), |
||||
make_tuple(ConversionTypes(COLOR_BGR2BGR565), 3, 2), |
||||
make_tuple(ConversionTypes(COLOR_RGBA2mRGBA), 4, 4), |
||||
make_tuple(ConversionTypes(COLOR_mRGBA2RGBA), 4, 4), |
||||
make_tuple(ConversionTypes(COLOR_YUV2RGB_NV12), 1, 3) |
||||
))) |
||||
{ |
||||
CvtColorParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const tuple<int, int, int> conversionParams = get<1>(params); |
||||
const int code = get<0>(conversionParams), scn = get<1>(conversionParams), |
||||
dcn = get<2>(conversionParams); |
||||
|
||||
UMat src(srcSize, CV_8UC(scn)), dst(srcSize, CV_8UC(scn)); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::cvtColor(src, dst, code, dcn); |
||||
|
||||
SANITY_CHECK(dst, 1); |
||||
} |
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
||||
@ -0,0 +1,302 @@ |
||||
/*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, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Fangfang Bai, fangfang@multicorewareinc.com
|
||||
// Jin Ma, jin@multicorewareinc.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*/
|
||||
|
||||
#include "perf_precomp.hpp" |
||||
#include "opencv2/ts/ocl_perf.hpp" |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
namespace cvtest { |
||||
namespace ocl { |
||||
|
||||
typedef tuple<Size, MatType, int> FilterParams; |
||||
typedef TestBaseWithParam<FilterParams> FilterFixture; |
||||
|
||||
///////////// Blur ////////////////////////
|
||||
|
||||
typedef FilterFixture BlurFixture; |
||||
|
||||
OCL_PERF_TEST_P(BlurFixture, Blur, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5))) |
||||
{ |
||||
const FilterParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), ksize = get<2>(params), bordertype = BORDER_CONSTANT; |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-5; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::blur(src, dst, Size(ksize, ksize), Point(-1, -1), bordertype); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
///////////// Laplacian////////////////////////
|
||||
|
||||
typedef FilterFixture LaplacianFixture; |
||||
|
||||
OCL_PERF_TEST_P(LaplacianFixture, Laplacian, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5))) |
||||
{ |
||||
const FilterParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), ksize = get<2>(params); |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-5; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::Laplacian(src, dst, -1, ksize, 1); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
///////////// Erode ////////////////////
|
||||
|
||||
typedef FilterFixture ErodeFixture; |
||||
|
||||
OCL_PERF_TEST_P(ErodeFixture, Erode, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5))) |
||||
{ |
||||
const FilterParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), ksize = get<2>(params); |
||||
const Mat ker = getStructuringElement(MORPH_RECT, Size(ksize, ksize)); |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst).in(ker); |
||||
|
||||
OCL_TEST_CYCLE() cv::erode(src, dst, ker); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
|
||||
///////////// Dilate ////////////////////
|
||||
|
||||
typedef FilterFixture DilateFixture; |
||||
|
||||
OCL_PERF_TEST_P(DilateFixture, Dilate, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5))) |
||||
{ |
||||
const FilterParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), ksize = get<2>(params); |
||||
const Mat ker = getStructuringElement(MORPH_RECT, Size(ksize, ksize)); |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst).in(ker); |
||||
|
||||
OCL_TEST_CYCLE() cv::dilate(src, dst, ker); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
|
||||
///////////// MorphologyEx ////////////////////////
|
||||
|
||||
CV_ENUM(MorphOp, MORPH_OPEN, MORPH_CLOSE, MORPH_GRADIENT, MORPH_TOPHAT, MORPH_BLACKHAT) |
||||
|
||||
typedef tuple<Size, MatType, MorphOp, int> MorphologyExParams; |
||||
typedef TestBaseWithParam<MorphologyExParams> MorphologyExFixture; |
||||
|
||||
OCL_PERF_TEST_P(MorphologyExFixture, MorphologyEx, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, MorphOp::all(), OCL_PERF_ENUM(3, 5))) |
||||
{ |
||||
const MorphologyExParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), op = get<2>(params), ksize = get<3>(params); |
||||
const Mat ker = getStructuringElement(MORPH_RECT, Size(ksize, ksize)); |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst).in(ker); |
||||
|
||||
OCL_TEST_CYCLE() cv::morphologyEx(src, dst, op, ker); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
|
||||
///////////// Sobel ////////////////////////
|
||||
|
||||
typedef Size_MatType SobelFixture; |
||||
|
||||
OCL_PERF_TEST_P(SobelFixture, Sobel, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) |
||||
{ |
||||
const Size_MatType_t params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), dx = 1, dy = 1; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type, sizeof(float) * 2); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::Sobel(src, dst, -1, dx, dy); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
|
||||
///////////// Scharr ////////////////////////
|
||||
|
||||
typedef Size_MatType ScharrFixture; |
||||
|
||||
OCL_PERF_TEST_P(ScharrFixture, Scharr, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) |
||||
{ |
||||
const Size_MatType_t params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), dx = 1, dy = 0; |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-5; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type, sizeof(float) * 2); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::Scharr(src, dst, -1, dx, dy); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
///////////// GaussianBlur ////////////////////////
|
||||
|
||||
typedef FilterFixture GaussianBlurFixture; |
||||
|
||||
OCL_PERF_TEST_P(GaussianBlurFixture, GaussianBlur, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5, 7))) |
||||
{ |
||||
const FilterParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), ksize = get<2>(params); |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 + DBL_EPSILON : 3e-4; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::GaussianBlur(src, dst, Size(ksize, ksize), 0); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
///////////// Filter2D ////////////////////////
|
||||
|
||||
typedef FilterFixture Filter2DFixture; |
||||
|
||||
OCL_PERF_TEST_P(Filter2DFixture, Filter2D, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, OCL_PERF_ENUM(3, 5))) |
||||
{ |
||||
const FilterParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), ksize = get<2>(params); |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-5; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type), kernel(ksize, ksize, CV_32SC1); |
||||
declare.in(src, WARMUP_RNG).in(kernel).out(dst); |
||||
randu(kernel, -3.0, 3.0); |
||||
|
||||
OCL_TEST_CYCLE() cv::filter2D(src, dst, -1, kernel); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
///////////// Bilateral ////////////////////////
|
||||
|
||||
typedef TestBaseWithParam<Size> BilateralFixture; |
||||
|
||||
OCL_PERF_TEST_P(BilateralFixture, Bilateral, OCL_TEST_SIZES) |
||||
{ |
||||
const Size srcSize = GetParam(); |
||||
const int d = 7; |
||||
const double sigmacolor = 50.0, sigmaspace = 50.0; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, CV_8UC1); |
||||
|
||||
UMat src(srcSize, CV_8UC1), dst(srcSize, CV_8UC1); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::bilateralFilter(src, dst, d, sigmacolor, sigmaspace); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
|
||||
///////////// MedianBlur ////////////////////////
|
||||
|
||||
typedef tuple<Size, int> MedianBlurParams; |
||||
typedef TestBaseWithParam<MedianBlurParams> MedianBlurFixture; |
||||
|
||||
OCL_PERF_TEST_P(MedianBlurFixture, Bilateral, ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(3, 5))) |
||||
{ |
||||
MedianBlurParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int ksize = get<1>(params); |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, CV_8UC1); |
||||
|
||||
UMat src(srcSize, CV_8UC1), dst(srcSize, CV_8UC1); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::medianBlur(src, dst, ksize); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
||||
@ -0,0 +1,246 @@ |
||||
/*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, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Fangfang Bai, fangfang@multicorewareinc.com
|
||||
// Jin Ma, jin@multicorewareinc.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*/
|
||||
|
||||
#include "perf_precomp.hpp" |
||||
#include "opencv2/ts/ocl_perf.hpp" |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
namespace cvtest { |
||||
namespace ocl { |
||||
|
||||
///////////// equalizeHist ////////////////////////
|
||||
|
||||
typedef TestBaseWithParam<Size> EqualizeHistFixture; |
||||
|
||||
OCL_PERF_TEST_P(EqualizeHistFixture, EqualizeHist, OCL_TEST_SIZES) |
||||
{ |
||||
const Size srcSize = GetParam(); |
||||
const double eps = 1; |
||||
|
||||
UMat src(srcSize, CV_8UC1), dst(srcSize, CV_8UC1); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::equalizeHist(src, dst); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
/////////// CopyMakeBorder //////////////////////
|
||||
|
||||
CV_ENUM(Border, BORDER_CONSTANT, BORDER_REPLICATE, BORDER_REFLECT, BORDER_WRAP, BORDER_REFLECT_101) |
||||
|
||||
typedef tuple<Size, MatType, Border> CopyMakeBorderParamType; |
||||
typedef TestBaseWithParam<CopyMakeBorderParamType> CopyMakeBorderFixture; |
||||
|
||||
OCL_PERF_TEST_P(CopyMakeBorderFixture, CopyMakeBorder, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, Border::all())) |
||||
{ |
||||
const CopyMakeBorderParamType params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), borderType = get<2>(params); |
||||
|
||||
UMat src(srcSize, type), dst; |
||||
const Size dstSize = srcSize + Size(12, 12); |
||||
dst.create(dstSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::copyMakeBorder(src, dst, 7, 5, 5, 7, borderType, cv::Scalar(1.0)); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
|
||||
///////////// CornerMinEigenVal ////////////////////////
|
||||
|
||||
typedef Size_MatType CornerMinEigenValFixture; |
||||
|
||||
OCL_PERF_TEST_P(CornerMinEigenValFixture, CornerMinEigenVal, |
||||
::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; |
||||
const int blockSize = 7, apertureSize = 1 + 2 * 3; |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, CV_32FC1); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
const int depth = CV_MAT_DEPTH(type); |
||||
const ERROR_TYPE errorType = depth == CV_8U ? ERROR_ABSOLUTE : ERROR_RELATIVE; |
||||
|
||||
OCL_TEST_CYCLE() cv::cornerMinEigenVal(src, dst, blockSize, apertureSize, borderType); |
||||
|
||||
SANITY_CHECK(dst, 1e-6, errorType); |
||||
} |
||||
|
||||
///////////// CornerHarris ////////////////////////
|
||||
|
||||
typedef Size_MatType CornerHarrisFixture; |
||||
|
||||
OCL_PERF_TEST_P(CornerHarrisFixture, CornerHarris, |
||||
::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::cornerHarris(src, dst, 5, 7, 0.1, borderType); |
||||
|
||||
SANITY_CHECK(dst, 5e-5); |
||||
} |
||||
|
||||
///////////// Integral ////////////////////////
|
||||
|
||||
typedef tuple<Size, MatDepth> IntegralParams; |
||||
typedef TestBaseWithParam<IntegralParams> IntegralFixture; |
||||
|
||||
OCL_PERF_TEST_P(IntegralFixture, Integral1, ::testing::Combine(OCL_TEST_SIZES, OCL_PERF_ENUM(CV_32S, CV_32F))) |
||||
{ |
||||
const IntegralParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int ddepth = get<1>(params); |
||||
|
||||
UMat src(srcSize, CV_8UC1), dst(srcSize + Size(1, 1), ddepth); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::integral(src, dst, ddepth); |
||||
|
||||
SANITY_CHECK(dst, 1e-6, ERROR_RELATIVE); |
||||
} |
||||
|
||||
///////////// Threshold ////////////////////////
|
||||
|
||||
CV_ENUM(ThreshType, THRESH_BINARY, THRESH_BINARY_INV, THRESH_TRUNC, THRESH_TOZERO_INV) |
||||
|
||||
typedef tuple<Size, MatType, ThreshType> ThreshParams; |
||||
typedef TestBaseWithParam<ThreshParams> ThreshFixture; |
||||
|
||||
OCL_PERF_TEST_P(ThreshFixture, Threshold, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, ThreshType::all())) |
||||
{ |
||||
const ThreshParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int srcType = get<1>(params); |
||||
const int threshType = get<2>(params); |
||||
const double maxValue = 220.0, threshold = 50; |
||||
|
||||
UMat src(srcSize, srcType), dst(srcSize, srcType); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::threshold(src, dst, threshold, maxValue, threshType); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
|
||||
///////////// CLAHE ////////////////////////
|
||||
|
||||
typedef TestBaseWithParam<Size> CLAHEFixture; |
||||
|
||||
OCL_PERF_TEST_P(CLAHEFixture, CLAHE, OCL_TEST_SIZES) |
||||
{ |
||||
const Size srcSize = GetParam(); |
||||
|
||||
UMat src(srcSize, CV_8UC1), dst(srcSize, CV_8UC1); |
||||
const double clipLimit = 40.0; |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
cv::Ptr<cv::CLAHE> clahe = cv::createCLAHE(clipLimit); |
||||
OCL_TEST_CYCLE() clahe->apply(src, dst); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
|
||||
///////////// SqrBoxFilter ////////////////////////
|
||||
|
||||
typedef TestBaseWithParam<Size> SqrBoxFilterFixture; |
||||
|
||||
OCL_PERF_TEST_P(SqrBoxFilterFixture, SqrBoxFilter, OCL_TEST_SIZES) |
||||
{ |
||||
const Size srcSize = GetParam(); |
||||
|
||||
UMat src(srcSize, CV_8UC1), dst(srcSize, CV_32SC1); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::sqrBoxFilter(src, dst, CV_32S, Size(3, 3)); |
||||
|
||||
SANITY_CHECK(dst); |
||||
} |
||||
|
||||
///////////// Canny ////////////////////////
|
||||
|
||||
typedef tuple<int, bool> CannyParams; |
||||
typedef TestBaseWithParam<CannyParams> CannyFixture; |
||||
|
||||
OCL_PERF_TEST_P(CannyFixture, Canny, ::testing::Combine(OCL_PERF_ENUM(3, 5), Bool())) |
||||
{ |
||||
const CannyParams params = GetParam(); |
||||
int apertureSize = get<0>(params); |
||||
bool L2Grad = get<1>(params); |
||||
|
||||
Mat _img = imread(getDataPath("gpu/stereobm/aloe-L.png"), cv::IMREAD_GRAYSCALE); |
||||
ASSERT_TRUE(!_img.empty()) << "can't open aloe-L.png"; |
||||
|
||||
UMat img; |
||||
_img.copyTo(img); |
||||
UMat edges(img.size(), CV_8UC1); |
||||
|
||||
declare.in(img, WARMUP_RNG).out(edges); |
||||
|
||||
OCL_TEST_CYCLE() cv::Canny(img, edges, 50.0, 100.0, apertureSize, L2Grad); |
||||
|
||||
if (apertureSize == 3) |
||||
SANITY_CHECK(edges); |
||||
else |
||||
SANITY_CHECK_NOTHING(); |
||||
} |
||||
|
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
||||
@ -0,0 +1,210 @@ |
||||
/*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, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Fangfang Bai, fangfang@multicorewareinc.com
|
||||
// Jin Ma, jin@multicorewareinc.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*/
|
||||
|
||||
#include "perf_precomp.hpp" |
||||
#include "opencv2/ts/ocl_perf.hpp" |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
namespace cvtest { |
||||
namespace ocl { |
||||
|
||||
///////////// WarpAffine ////////////////////////
|
||||
|
||||
CV_ENUM(InterType, INTER_NEAREST, INTER_LINEAR) |
||||
|
||||
typedef tuple<Size, MatType, InterType> WarpAffineParams; |
||||
typedef TestBaseWithParam<WarpAffineParams> WarpAffineFixture; |
||||
|
||||
OCL_PERF_TEST_P(WarpAffineFixture, WarpAffine, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, InterType::all())) |
||||
{ |
||||
static const double coeffs[2][3] = |
||||
{ |
||||
{ cos(CV_PI / 6), -sin(CV_PI / 6), 100.0 }, |
||||
{ sin(CV_PI / 6), cos(CV_PI / 6) , -100.0 } |
||||
}; |
||||
Mat M(2, 3, CV_64F, (void *)coeffs); |
||||
|
||||
const WarpAffineParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), interpolation = get<2>(params); |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-4; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::warpAffine(src, dst, M, srcSize, interpolation); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
///////////// WarpPerspective ////////////////////////
|
||||
|
||||
typedef WarpAffineParams WarpPerspectiveParams; |
||||
typedef TestBaseWithParam<WarpPerspectiveParams> WarpPerspectiveFixture; |
||||
|
||||
OCL_PERF_TEST_P(WarpPerspectiveFixture, WarpPerspective, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, InterType::all())) |
||||
{ |
||||
static const double coeffs[3][3] = |
||||
{ |
||||
{cos(CV_PI / 6), -sin(CV_PI / 6), 100.0}, |
||||
{sin(CV_PI / 6), cos(CV_PI / 6), -100.0}, |
||||
{0.0, 0.0, 1.0} |
||||
}; |
||||
Mat M(3, 3, CV_64F, (void *)coeffs); |
||||
|
||||
const WarpPerspectiveParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), interpolation = get<2>(params); |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-4; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::warpPerspective(src, dst, M, srcSize, interpolation); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
///////////// Resize ////////////////////////
|
||||
|
||||
typedef tuple<Size, MatType, InterType, double> ResizeParams; |
||||
typedef TestBaseWithParam<ResizeParams> ResizeFixture; |
||||
|
||||
OCL_PERF_TEST_P(ResizeFixture, Resize, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, |
||||
InterType::all(), ::testing::Values(0.5, 2.0))) |
||||
{ |
||||
const ResizeParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), interType = get<2>(params); |
||||
double scale = get<3>(params); |
||||
const Size dstSize(cvRound(srcSize.width * scale), cvRound(srcSize.height * scale)); |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-4; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
checkDeviceMaxMemoryAllocSize(dstSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(dstSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::resize(src, dst, Size(), scale, scale, interType); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
typedef tuple<Size, MatType, double> ResizeAreaParams; |
||||
typedef TestBaseWithParam<ResizeAreaParams> ResizeAreaFixture; |
||||
|
||||
OCL_PERF_TEST_P(ResizeAreaFixture, Resize, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, ::testing::Values(0.3, 0.5, 0.6))) |
||||
{ |
||||
const ResizeAreaParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params); |
||||
double scale = get<2>(params); |
||||
const Size dstSize(cvRound(srcSize.width * scale), cvRound(srcSize.height * scale)); |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-4; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
checkDeviceMaxMemoryAllocSize(dstSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(dstSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::resize(src, dst, Size(), scale, scale, cv::INTER_AREA); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
///////////// Remap ////////////////////////
|
||||
|
||||
typedef tuple<Size, MatType, InterType> RemapParams; |
||||
typedef TestBaseWithParam<RemapParams> RemapFixture; |
||||
|
||||
OCL_PERF_TEST_P(RemapFixture, Remap, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES, InterType::all())) |
||||
{ |
||||
const RemapParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params), interpolation = get<2>(params), borderMode = BORDER_CONSTANT; |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-4; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(srcSize, type); |
||||
UMat xmap(srcSize, CV_32FC1), ymap(srcSize, CV_32FC1); |
||||
|
||||
{ |
||||
Mat _xmap = xmap.getMat(ACCESS_WRITE), _ymap = ymap.getMat(ACCESS_WRITE); |
||||
for (int i = 0; i < srcSize.height; ++i) |
||||
{ |
||||
float * const xmap_row = _xmap.ptr<float>(i); |
||||
float * const ymap_row = _ymap.ptr<float>(i); |
||||
|
||||
for (int j = 0; j < srcSize.width; ++j) |
||||
{ |
||||
xmap_row[j] = (j - srcSize.width * 0.5f) * 0.75f + srcSize.width * 0.5f; |
||||
ymap_row[j] = (i - srcSize.height * 0.5f) * 0.75f + srcSize.height * 0.5f; |
||||
} |
||||
} |
||||
} |
||||
declare.in(src, WARMUP_RNG).in(xmap, ymap, WARMUP_READ).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::remap(src, dst, xmap, ymap, interpolation, borderMode); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
||||
@ -0,0 +1,78 @@ |
||||
/*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, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Fangfang Bai, fangfang@multicorewareinc.com
|
||||
// Jin Ma, jin@multicorewareinc.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*/
|
||||
|
||||
#include "perf_precomp.hpp" |
||||
#include "opencv2/ts/ocl_perf.hpp" |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
namespace cvtest { |
||||
namespace ocl { |
||||
|
||||
///////////// Moments ////////////////////////
|
||||
|
||||
typedef tuple<Size, bool> MomentsParams; |
||||
typedef TestBaseWithParam<MomentsParams> MomentsFixture; |
||||
|
||||
OCL_PERF_TEST_P(MomentsFixture, Moments, |
||||
::testing::Combine(OCL_TEST_SIZES, ::testing::Bool())) |
||||
{ |
||||
const MomentsParams params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const bool binaryImage = get<1>(params); |
||||
|
||||
cv::Moments m; |
||||
UMat src(srcSize, CV_8UC1); |
||||
declare.in(src, WARMUP_RNG); |
||||
|
||||
OCL_TEST_CYCLE() m = cv::moments(src, binaryImage); |
||||
|
||||
SANITY_CHECK_MOMENTS(m, 1e-6, ERROR_RELATIVE); |
||||
} |
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
||||
@ -0,0 +1,105 @@ |
||||
/*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, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Fangfang Bai, fangfang@multicorewareinc.com
|
||||
// Jin Ma, jin@multicorewareinc.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*/
|
||||
|
||||
#include "perf_precomp.hpp" |
||||
#include "opencv2/ts/ocl_perf.hpp" |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
namespace cvtest { |
||||
namespace ocl { |
||||
|
||||
///////////// PyrDown //////////////////////
|
||||
|
||||
typedef Size_MatType PyrDownFixture; |
||||
|
||||
OCL_PERF_TEST_P(PyrDownFixture, PyrDown, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) |
||||
{ |
||||
const Size_MatType_t params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params); |
||||
const Size dstSize((srcSize.height + 1) >> 1, (srcSize.width + 1) >> 1); |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-5; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
checkDeviceMaxMemoryAllocSize(dstSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(dstSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::pyrDown(src, dst); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
///////////// PyrUp ////////////////////////
|
||||
|
||||
typedef Size_MatType PyrUpFixture; |
||||
|
||||
OCL_PERF_TEST_P(PyrUpFixture, PyrUp, |
||||
::testing::Combine(OCL_TEST_SIZES, OCL_TEST_TYPES)) |
||||
{ |
||||
const Size_MatType_t params = GetParam(); |
||||
const Size srcSize = get<0>(params); |
||||
const int type = get<1>(params); |
||||
const Size dstSize(srcSize.height << 1, srcSize.width << 1); |
||||
const double eps = CV_MAT_DEPTH(type) <= CV_32S ? 1 : 1e-5; |
||||
|
||||
checkDeviceMaxMemoryAllocSize(srcSize, type); |
||||
checkDeviceMaxMemoryAllocSize(dstSize, type); |
||||
|
||||
UMat src(srcSize, type), dst(dstSize, type); |
||||
declare.in(src, WARMUP_RNG).out(dst); |
||||
|
||||
OCL_TEST_CYCLE() cv::pyrDown(src, dst); |
||||
|
||||
SANITY_CHECK(dst, eps); |
||||
} |
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
||||
@ -0,0 +1,514 @@ |
||||
/*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, Multicoreware, Inc., all rights reserved. |
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved. |
||||
// Third party copyrights are property of their respective owners. |
||||
// |
||||
// @Authors |
||||
// Peng Xiao, pengxiao@multicorewareinc.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*/ |
||||
|
||||
// Smoothing perpendicular to the derivative direction with a triangle filter |
||||
// only support 3x3 Sobel kernel |
||||
// h (-1) = 1, h (0) = 2, h (1) = 1 |
||||
// h'(-1) = -1, h'(0) = 0, h'(1) = 1 |
||||
// thus sobel 2D operator can be calculated as: |
||||
// h'(x, y) = h'(x)h(y) for x direction |
||||
// |
||||
// src input 8bit single channel image data |
||||
// dx_buf output dx buffer |
||||
// dy_buf output dy buffer |
||||
|
||||
__kernel void __attribute__((reqd_work_group_size(16, 16, 1))) |
||||
calcSobelRowPass |
||||
(__global const uchar * src, int src_step, int src_offset, int rows, int cols, |
||||
__global uchar * dx_buf, int dx_buf_step, int dx_buf_offset, |
||||
__global uchar * dy_buf, int dy_buf_step, int dy_buf_offset) |
||||
{ |
||||
int gidx = get_global_id(0); |
||||
int gidy = get_global_id(1); |
||||
|
||||
int lidx = get_local_id(0); |
||||
int lidy = get_local_id(1); |
||||
|
||||
__local int smem[16][18]; |
||||
|
||||
smem[lidy][lidx + 1] = src[mad24(src_step, min(gidy, rows - 1), gidx + src_offset)]; |
||||
if (lidx == 0) |
||||
{ |
||||
smem[lidy][0] = src[mad24(src_step, min(gidy, rows - 1), max(gidx - 1, 0) + src_offset)]; |
||||
smem[lidy][17] = src[mad24(src_step, min(gidy, rows - 1), min(gidx + 16, cols - 1) + src_offset)]; |
||||
} |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (gidy < rows && gidx < cols) |
||||
{ |
||||
*(__global short *)(dx_buf + mad24(gidy, dx_buf_step, gidx * (int)sizeof(short) + dx_buf_offset)) = |
||||
smem[lidy][lidx + 2] - smem[lidy][lidx]; |
||||
*(__global short *)(dy_buf + mad24(gidy, dy_buf_step, gidx * (int)sizeof(short) + dy_buf_offset)) = |
||||
smem[lidy][lidx] + 2 * smem[lidy][lidx + 1] + smem[lidy][lidx + 2]; |
||||
} |
||||
} |
||||
|
||||
inline int calc(short x, short y) |
||||
{ |
||||
#ifdef L2GRAD |
||||
return x * x + y * y; |
||||
#else |
||||
return (x >= 0 ? x : -x) + (y >= 0 ? y : -y); |
||||
#endif |
||||
} |
||||
|
||||
// calculate the magnitude of the filter pass combining both x and y directions |
||||
// This is the non-buffered version(non-3x3 sobel) |
||||
// |
||||
// dx_buf dx buffer, calculated from calcSobelRowPass |
||||
// dy_buf dy buffer, calculated from calcSobelRowPass |
||||
// dx direvitive in x direction output |
||||
// dy direvitive in y direction output |
||||
// mag magnitude direvitive of xy output |
||||
|
||||
__kernel void calcMagnitude(__global const uchar * dxptr, int dx_step, int dx_offset, |
||||
__global const uchar * dyptr, int dy_step, int dy_offset, |
||||
__global uchar * magptr, int mag_step, int mag_offset, int rows, int cols) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (y < rows && x < cols) |
||||
{ |
||||
int dx_index = mad24(dx_step, y, x * (int)sizeof(short) + dx_offset); |
||||
int dy_index = mad24(dy_step, y, x * (int)sizeof(short) + dy_offset); |
||||
int mag_index = mad24(mag_step, y + 1, (x + 1) * (int)sizeof(int) + mag_offset); |
||||
|
||||
__global const short * dx = (__global const short *)(dxptr + dx_index); |
||||
__global const short * dy = (__global const short *)(dyptr + dy_index); |
||||
__global int * mag = (__global int *)(magptr + mag_index); |
||||
|
||||
mag[0] = calc(dx[0], dy[0]); |
||||
} |
||||
} |
||||
|
||||
// calculate the magnitude of the filter pass combining both x and y directions |
||||
// This is the buffered version(3x3 sobel) |
||||
// |
||||
// dx_buf dx buffer, calculated from calcSobelRowPass |
||||
// dy_buf dy buffer, calculated from calcSobelRowPass |
||||
// dx direvitive in x direction output |
||||
// dy direvitive in y direction output |
||||
// mag magnitude direvitive of xy output |
||||
__kernel void __attribute__((reqd_work_group_size(16, 16, 1))) |
||||
calcMagnitude_buf |
||||
(__global const short * dx_buf, int dx_buf_step, int dx_buf_offset, |
||||
__global const short * dy_buf, int dy_buf_step, int dy_buf_offset, |
||||
__global short * dx, int dx_step, int dx_offset, |
||||
__global short * dy, int dy_step, int dy_offset, |
||||
__global int * mag, int mag_step, int mag_offset, |
||||
int rows, int cols) |
||||
{ |
||||
dx_buf_step /= sizeof(*dx_buf); |
||||
dx_buf_offset /= sizeof(*dx_buf); |
||||
dy_buf_step /= sizeof(*dy_buf); |
||||
dy_buf_offset /= sizeof(*dy_buf); |
||||
dx_step /= sizeof(*dx); |
||||
dx_offset /= sizeof(*dx); |
||||
dy_step /= sizeof(*dy); |
||||
dy_offset /= sizeof(*dy); |
||||
mag_step /= sizeof(*mag); |
||||
mag_offset /= sizeof(*mag); |
||||
|
||||
int gidx = get_global_id(0); |
||||
int gidy = get_global_id(1); |
||||
|
||||
int lidx = get_local_id(0); |
||||
int lidy = get_local_id(1); |
||||
|
||||
__local short sdx[18][16]; |
||||
__local short sdy[18][16]; |
||||
|
||||
sdx[lidy + 1][lidx] = dx_buf[gidx + min(gidy, rows - 1) * dx_buf_step + dx_buf_offset]; |
||||
sdy[lidy + 1][lidx] = dy_buf[gidx + min(gidy, rows - 1) * dy_buf_step + dy_buf_offset]; |
||||
if (lidy == 0) |
||||
{ |
||||
sdx[0][lidx] = dx_buf[gidx + min(max(gidy - 1, 0), rows - 1) * dx_buf_step + dx_buf_offset]; |
||||
sdx[17][lidx] = dx_buf[gidx + min(gidy + 16, rows - 1) * dx_buf_step + dx_buf_offset]; |
||||
|
||||
sdy[0][lidx] = dy_buf[gidx + min(max(gidy - 1, 0), rows - 1) * dy_buf_step + dy_buf_offset]; |
||||
sdy[17][lidx] = dy_buf[gidx + min(gidy + 16, rows - 1) * dy_buf_step + dy_buf_offset]; |
||||
} |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (gidx < cols && gidy < rows) |
||||
{ |
||||
short x = sdx[lidy][lidx] + 2 * sdx[lidy + 1][lidx] + sdx[lidy + 2][lidx]; |
||||
short y = -sdy[lidy][lidx] + sdy[lidy + 2][lidx]; |
||||
|
||||
dx[gidx + gidy * dx_step + dx_offset] = x; |
||||
dy[gidx + gidy * dy_step + dy_offset] = y; |
||||
|
||||
mag[(gidx + 1) + (gidy + 1) * mag_step + mag_offset] = calc(x, y); |
||||
} |
||||
} |
||||
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////////////// |
||||
// 0.4142135623730950488016887242097 is tan(22.5) |
||||
|
||||
#define CANNY_SHIFT 15 |
||||
#define TG22 (int)(0.4142135623730950488016887242097f*(1<<CANNY_SHIFT) + 0.5f) |
||||
|
||||
// First pass of edge detection and non-maximum suppression |
||||
// edgetype is set to for each pixel: |
||||
// 0 - below low thres, not an edge |
||||
// 1 - maybe an edge |
||||
// 2 - is an edge, either magnitude is greater than high thres, or |
||||
// Given estimates of the image gradients, a search is then carried out |
||||
// to determine if the gradient magnitude assumes a local maximum in the gradient direction. |
||||
// if the rounded gradient angle is zero degrees (i.e. the edge is in the north-south direction) the point will be considered to be on the edge if its gradient magnitude is greater than the magnitudes in the west and east directions, |
||||
// if the rounded gradient angle is 90 degrees (i.e. the edge is in the east-west direction) the point will be considered to be on the edge if its gradient magnitude is greater than the magnitudes in the north and south directions, |
||||
// if the rounded gradient angle is 135 degrees (i.e. the edge is in the north east-south west direction) the point will be considered to be on the edge if its gradient magnitude is greater than the magnitudes in the north west and south east directions, |
||||
// if the rounded gradient angle is 45 degrees (i.e. the edge is in the north west-south east direction)the point will be considered to be on the edge if its gradient magnitude is greater than the magnitudes in the north east and south west directions. |
||||
// |
||||
// dx, dy direvitives of x and y direction |
||||
// mag magnitudes calculated from calcMagnitude function |
||||
// map output containing raw edge types |
||||
|
||||
__kernel void __attribute__((reqd_work_group_size(16,16,1))) |
||||
calcMap( |
||||
__global const uchar * dx, int dx_step, int dx_offset, |
||||
__global const uchar * dy, int dy_step, int dy_offset, |
||||
__global const uchar * mag, int mag_step, int mag_offset, |
||||
__global uchar * map, int map_step, int map_offset, |
||||
int rows, int cols, int low_thresh, int high_thresh) |
||||
{ |
||||
__local int smem[18][18]; |
||||
|
||||
int gidx = get_global_id(0); |
||||
int gidy = get_global_id(1); |
||||
|
||||
int lidx = get_local_id(0); |
||||
int lidy = get_local_id(1); |
||||
|
||||
int grp_idx = get_global_id(0) & 0xFFFFF0; |
||||
int grp_idy = get_global_id(1) & 0xFFFFF0; |
||||
|
||||
int tid = lidx + lidy * 16; |
||||
int lx = tid % 18; |
||||
int ly = tid / 18; |
||||
|
||||
mag += mag_offset; |
||||
if (ly < 14) |
||||
smem[ly][lx] = *(__global const int *)(mag + |
||||
mad24(mag_step, min(grp_idy + ly, rows - 1), (int)sizeof(int) * (grp_idx + lx))); |
||||
if (ly < 4 && grp_idy + ly + 14 <= rows && grp_idx + lx <= cols) |
||||
smem[ly + 14][lx] = *(__global const int *)(mag + |
||||
mad24(mag_step, min(grp_idy + ly + 14, rows - 1), (int)sizeof(int) * (grp_idx + lx))); |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (gidy < rows && gidx < cols) |
||||
{ |
||||
// 0 - the pixel can not belong to an edge |
||||
// 1 - the pixel might belong to an edge |
||||
// 2 - the pixel does belong to an edge |
||||
int edge_type = 0; |
||||
int m = smem[lidy + 1][lidx + 1]; |
||||
|
||||
if (m > low_thresh) |
||||
{ |
||||
short xs = *(__global const short *)(dx + mad24(gidy, dx_step, dx_offset + (int)sizeof(short) * gidx)); |
||||
short ys = *(__global const short *)(dy + mad24(gidy, dy_step, dy_offset + (int)sizeof(short) * gidx)); |
||||
int x = abs(xs), y = abs(ys); |
||||
|
||||
int tg22x = x * TG22; |
||||
y <<= CANNY_SHIFT; |
||||
|
||||
if (y < tg22x) |
||||
{ |
||||
if (m > smem[lidy + 1][lidx] && m >= smem[lidy + 1][lidx + 2]) |
||||
edge_type = 1 + (int)(m > high_thresh); |
||||
} |
||||
else |
||||
{ |
||||
int tg67x = tg22x + (x << (1 + CANNY_SHIFT)); |
||||
if (y > tg67x) |
||||
{ |
||||
if (m > smem[lidy][lidx + 1]&& m >= smem[lidy + 2][lidx + 1]) |
||||
edge_type = 1 + (int)(m > high_thresh); |
||||
} |
||||
else |
||||
{ |
||||
int s = (xs ^ ys) < 0 ? -1 : 1; |
||||
if (m > smem[lidy][lidx + 1 - s]&& m > smem[lidy + 2][lidx + 1 + s]) |
||||
edge_type = 1 + (int)(m > high_thresh); |
||||
} |
||||
} |
||||
} |
||||
*(__global int *)(map + mad24(map_step, gidy + 1, (gidx + 1) * (int)sizeof(int) + map_offset)) = edge_type; |
||||
} |
||||
} |
||||
|
||||
#undef CANNY_SHIFT |
||||
#undef TG22 |
||||
|
||||
struct PtrStepSz |
||||
{ |
||||
__global uchar * ptr; |
||||
int step, rows, cols; |
||||
}; |
||||
|
||||
inline int get(struct PtrStepSz data, int y, int x) |
||||
{ |
||||
return *(__global int *)(data.ptr + mad24(data.step, y + 1, (int)sizeof(int) * (x + 1))); |
||||
} |
||||
|
||||
inline void set(struct PtrStepSz data, int y, int x, int value) |
||||
{ |
||||
*(__global int *)(data.ptr + mad24(data.step, y + 1, (int)sizeof(int) * (x + 1))) = value; |
||||
} |
||||
|
||||
// perform Hysteresis for pixel whose edge type is 1 |
||||
// |
||||
// If candidate pixel (edge type is 1) has a neighbour pixel (in 3x3 area) with type 2, it is believed to be part of an edge and |
||||
// marked as edge. Each thread will iterate for 16 times to connect local edges. |
||||
// Candidate pixel being identified as edge will then be tested if there is nearby potiential edge points. If there is, counter will |
||||
// be incremented by 1 and the point location is stored. These potiential candidates will be processed further in next kernel. |
||||
// |
||||
// map raw edge type results calculated from calcMap. |
||||
// stack the potiential edge points found in this kernel call |
||||
// counter the number of potiential edge points |
||||
|
||||
__kernel void __attribute__((reqd_work_group_size(16,16,1))) |
||||
edgesHysteresisLocal |
||||
(__global uchar * map_ptr, int map_step, int map_offset, |
||||
__global ushort2 * st, __global unsigned int * counter, |
||||
int rows, int cols) |
||||
{ |
||||
struct PtrStepSz map = { map_ptr + map_offset, map_step, rows + 1, cols + 1 }; |
||||
|
||||
__local int smem[18][18]; |
||||
|
||||
int2 blockIdx = (int2)(get_group_id(0), get_group_id(1)); |
||||
int2 blockDim = (int2)(get_local_size(0), get_local_size(1)); |
||||
int2 threadIdx = (int2)(get_local_id(0), get_local_id(1)); |
||||
|
||||
const int x = blockIdx.x * blockDim.x + threadIdx.x; |
||||
const int y = blockIdx.y * blockDim.y + threadIdx.y; |
||||
|
||||
smem[threadIdx.y + 1][threadIdx.x + 1] = x < map.cols && y < map.rows ? get(map, y, x) : 0; |
||||
if (threadIdx.y == 0) |
||||
smem[0][threadIdx.x + 1] = x < map.cols ? get(map, y - 1, x) : 0; |
||||
if (threadIdx.y == blockDim.y - 1) |
||||
smem[blockDim.y + 1][threadIdx.x + 1] = y + 1 < map.rows ? get(map, y + 1, x) : 0; |
||||
if (threadIdx.x == 0) |
||||
smem[threadIdx.y + 1][0] = y < map.rows ? get(map, y, x - 1) : 0; |
||||
if (threadIdx.x == blockDim.x - 1) |
||||
smem[threadIdx.y + 1][blockDim.x + 1] = x + 1 < map.cols && y < map.rows ? get(map, y, x + 1) : 0; |
||||
if (threadIdx.x == 0 && threadIdx.y == 0) |
||||
smem[0][0] = y > 0 && x > 0 ? get(map, y - 1, x - 1) : 0; |
||||
if (threadIdx.x == blockDim.x - 1 && threadIdx.y == 0) |
||||
smem[0][blockDim.x + 1] = y > 0 && x + 1 < map.cols ? get(map, y - 1, x + 1) : 0; |
||||
if (threadIdx.x == 0 && threadIdx.y == blockDim.y - 1) |
||||
smem[blockDim.y + 1][0] = y + 1 < map.rows && x > 0 ? get(map, y + 1, x - 1) : 0; |
||||
if (threadIdx.x == blockDim.x - 1 && threadIdx.y == blockDim.y - 1) |
||||
smem[blockDim.y + 1][blockDim.x + 1] = y + 1 < map.rows && x + 1 < map.cols ? get(map, y + 1, x + 1) : 0; |
||||
|
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (x >= cols || y >= rows) |
||||
return; |
||||
|
||||
int n; |
||||
|
||||
#pragma unroll |
||||
for (int k = 0; k < 16; ++k) |
||||
{ |
||||
n = 0; |
||||
|
||||
if (smem[threadIdx.y + 1][threadIdx.x + 1] == 1) |
||||
{ |
||||
n += smem[threadIdx.y ][threadIdx.x ] == 2; |
||||
n += smem[threadIdx.y ][threadIdx.x + 1] == 2; |
||||
n += smem[threadIdx.y ][threadIdx.x + 2] == 2; |
||||
|
||||
n += smem[threadIdx.y + 1][threadIdx.x ] == 2; |
||||
n += smem[threadIdx.y + 1][threadIdx.x + 2] == 2; |
||||
|
||||
n += smem[threadIdx.y + 2][threadIdx.x ] == 2; |
||||
n += smem[threadIdx.y + 2][threadIdx.x + 1] == 2; |
||||
n += smem[threadIdx.y + 2][threadIdx.x + 2] == 2; |
||||
} |
||||
|
||||
if (n > 0) |
||||
smem[threadIdx.y + 1][threadIdx.x + 1] = 2; |
||||
} |
||||
|
||||
const int e = smem[threadIdx.y + 1][threadIdx.x + 1]; |
||||
set(map, y, x, e); |
||||
n = 0; |
||||
|
||||
if (e == 2) |
||||
{ |
||||
n += smem[threadIdx.y ][threadIdx.x ] == 1; |
||||
n += smem[threadIdx.y ][threadIdx.x + 1] == 1; |
||||
n += smem[threadIdx.y ][threadIdx.x + 2] == 1; |
||||
|
||||
n += smem[threadIdx.y + 1][threadIdx.x ] == 1; |
||||
n += smem[threadIdx.y + 1][threadIdx.x + 2] == 1; |
||||
|
||||
n += smem[threadIdx.y + 2][threadIdx.x ] == 1; |
||||
n += smem[threadIdx.y + 2][threadIdx.x + 1] == 1; |
||||
n += smem[threadIdx.y + 2][threadIdx.x + 2] == 1; |
||||
} |
||||
|
||||
if (n > 0) |
||||
{ |
||||
const int ind = atomic_inc(counter); |
||||
st[ind] = (ushort2)(x + 1, y + 1); |
||||
} |
||||
} |
||||
|
||||
__constant int c_dx[8] = {-1, 0, 1, -1, 1, -1, 0, 1}; |
||||
__constant int c_dy[8] = {-1, -1, -1, 0, 0, 1, 1, 1}; |
||||
|
||||
|
||||
#define stack_size 512 |
||||
#define map_index mad24(map_step, pos.y, pos.x * (int)sizeof(int)) |
||||
|
||||
__kernel void __attribute__((reqd_work_group_size(128, 1, 1))) |
||||
edgesHysteresisGlobal(__global uchar * map, int map_step, int map_offset, |
||||
__global ushort2 * st1, __global ushort2 * st2, __global int * counter, |
||||
int rows, int cols, int count) |
||||
{ |
||||
map += map_offset; |
||||
|
||||
int lidx = get_local_id(0); |
||||
|
||||
int grp_idx = get_group_id(0); |
||||
int grp_idy = get_group_id(1); |
||||
|
||||
__local unsigned int s_counter, s_ind; |
||||
__local ushort2 s_st[stack_size]; |
||||
|
||||
if (lidx == 0) |
||||
s_counter = 0; |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
int ind = mad24(grp_idy, (int)get_local_size(0), grp_idx); |
||||
|
||||
if (ind < count) |
||||
{ |
||||
ushort2 pos = st1[ind]; |
||||
if (lidx < 8) |
||||
{ |
||||
pos.x += c_dx[lidx]; |
||||
pos.y += c_dy[lidx]; |
||||
if (pos.x > 0 && pos.x <= cols && pos.y > 0 && pos.y <= rows && *(__global int *)(map + map_index) == 1) |
||||
{ |
||||
*(__global int *)(map + map_index) = 2; |
||||
ind = atomic_inc(&s_counter); |
||||
s_st[ind] = pos; |
||||
} |
||||
} |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
while (s_counter > 0 && s_counter <= stack_size - get_local_size(0)) |
||||
{ |
||||
const int subTaskIdx = lidx >> 3; |
||||
const int portion = min(s_counter, (uint)(get_local_size(0)>> 3)); |
||||
|
||||
if (subTaskIdx < portion) |
||||
pos = s_st[s_counter - 1 - subTaskIdx]; |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (lidx == 0) |
||||
s_counter -= portion; |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
if (subTaskIdx < portion) |
||||
{ |
||||
pos.x += c_dx[lidx & 7]; |
||||
pos.y += c_dy[lidx & 7]; |
||||
if (pos.x > 0 && pos.x <= cols && pos.y > 0 && pos.y <= rows && *(__global int *)(map + map_index) == 1) |
||||
{ |
||||
*(__global int *)(map + map_index) = 2; |
||||
ind = atomic_inc(&s_counter); |
||||
s_st[ind] = pos; |
||||
} |
||||
} |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
} |
||||
|
||||
if (s_counter > 0) |
||||
{ |
||||
if (lidx == 0) |
||||
{ |
||||
ind = atomic_add(counter, s_counter); |
||||
s_ind = ind - s_counter; |
||||
} |
||||
barrier(CLK_LOCAL_MEM_FENCE); |
||||
|
||||
ind = s_ind; |
||||
for (int i = lidx; i < (int)s_counter; i += get_local_size(0)) |
||||
st2[ind + i] = s_st[i]; |
||||
} |
||||
} |
||||
} |
||||
|
||||
#undef map_index |
||||
#undef stack_size |
||||
|
||||
// Get the edge result. egde type of value 2 will be marked as an edge point and set to 255. Otherwise 0. |
||||
// map edge type mappings |
||||
// dst edge output |
||||
|
||||
__kernel void getEdges(__global const uchar * mapptr, int map_step, int map_offset, |
||||
__global uchar * dst, int dst_step, int dst_offset, int rows, int cols) |
||||
{ |
||||
int x = get_global_id(0); |
||||
int y = get_global_id(1); |
||||
|
||||
if (y < rows && x < cols) |
||||
{ |
||||
int map_index = mad24(map_step, y + 1, (x + 1) * (int)sizeof(int) + map_offset); |
||||
int dst_index = mad24(dst_step, y, x + dst_offset); |
||||
|
||||
__global const int * map = (__global const int *)(mapptr + map_index); |
||||
|
||||
dst[dst_index] = (uchar)(-(map[0] >> 1)); |
||||
} |
||||
} |
||||
@ -0,0 +1,117 @@ |
||||
/*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, Multicoreware, Inc., all rights reserved.
|
||||
// Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
|
||||
// Third party copyrights are property of their respective owners.
|
||||
//
|
||||
// @Authors
|
||||
// Peng Xiao, pengxiao@multicorewareinc.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*/
|
||||
|
||||
#include "test_precomp.hpp" |
||||
#include "opencv2/ts/ocl_test.hpp" |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
namespace cvtest { |
||||
namespace ocl { |
||||
|
||||
////////////////////////////////////////////////////////
|
||||
// Canny
|
||||
|
||||
IMPLEMENT_PARAM_CLASS(AppertureSize, int) |
||||
IMPLEMENT_PARAM_CLASS(L2gradient, bool) |
||||
IMPLEMENT_PARAM_CLASS(UseRoi, bool) |
||||
|
||||
PARAM_TEST_CASE(Canny, AppertureSize, L2gradient, UseRoi) |
||||
{ |
||||
int apperture_size; |
||||
bool useL2gradient, use_roi; |
||||
|
||||
TEST_DECLARE_INPUT_PARAMETER(src) |
||||
TEST_DECLARE_OUTPUT_PARAMETER(dst) |
||||
|
||||
virtual void SetUp() |
||||
{ |
||||
apperture_size = GET_PARAM(0); |
||||
useL2gradient = GET_PARAM(1); |
||||
use_roi = GET_PARAM(2); |
||||
} |
||||
|
||||
void generateTestData() |
||||
{ |
||||
Mat img = readImage("shared/fruits.png", IMREAD_GRAYSCALE); |
||||
ASSERT_FALSE(img.empty()) << "cann't load shared/fruits.png"; |
||||
|
||||
Size roiSize = img.size(); |
||||
int type = img.type(); |
||||
ASSERT_EQ(CV_8UC1, type); |
||||
|
||||
Border srcBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); |
||||
randomSubMat(src, src_roi, roiSize, srcBorder, type, 2, 100); |
||||
img.copyTo(src_roi); |
||||
|
||||
Border dstBorder = randomBorder(0, use_roi ? MAX_VALUE : 0); |
||||
randomSubMat(dst, dst_roi, roiSize, dstBorder, type, 5, 16); |
||||
|
||||
UMAT_UPLOAD_INPUT_PARAMETER(src) |
||||
UMAT_UPLOAD_OUTPUT_PARAMETER(dst) |
||||
} |
||||
}; |
||||
|
||||
OCL_TEST_P(Canny, Accuracy) |
||||
{ |
||||
generateTestData(); |
||||
|
||||
const double low_thresh = 50.0, high_thresh = 100.0; |
||||
|
||||
OCL_OFF(cv::Canny(src_roi, dst_roi, low_thresh, high_thresh, apperture_size, useL2gradient)); |
||||
OCL_ON(cv::Canny(usrc_roi, udst_roi, low_thresh, high_thresh, apperture_size, useL2gradient)); |
||||
|
||||
EXPECT_MAT_SIMILAR(dst_roi, udst_roi, 1e-2); |
||||
EXPECT_MAT_SIMILAR(dst, udst, 1e-2); |
||||
} |
||||
|
||||
OCL_INSTANTIATE_TEST_CASE_P(ImgProc, Canny, testing::Combine( |
||||
testing::Values(AppertureSize(3), AppertureSize(5)), |
||||
testing::Values(L2gradient(false), L2gradient(true)), |
||||
testing::Values(UseRoi(false), UseRoi(true)))); |
||||
|
||||
} } // namespace cvtest::ocl
|
||||
|
||||
#endif // HAVE_OPENCL
|
||||
@ -0,0 +1,61 @@ |
||||
#include "perf_precomp.hpp" |
||||
#include <opencv2/imgproc.hpp> |
||||
|
||||
#include "opencv2/ts/ocl_perf.hpp" |
||||
|
||||
using namespace std; |
||||
using namespace cv; |
||||
using namespace perf; |
||||
using std::tr1::make_tuple; |
||||
using std::tr1::get; |
||||
|
||||
typedef std::tr1::tuple<std::string, std::string, int> Cascade_Image_MinSize_t; |
||||
typedef perf::TestBaseWithParam<Cascade_Image_MinSize_t> Cascade_Image_MinSize; |
||||
|
||||
#ifdef HAVE_OPENCL |
||||
|
||||
OCL_PERF_TEST_P(Cascade_Image_MinSize, CascadeClassifier, |
||||
testing::Combine( |
||||
testing::Values( string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt.xml"), |
||||
string("cv/cascadeandhog/cascades/haarcascade_frontalface_alt_old.xml"), |
||||
string("cv/cascadeandhog/cascades/lbpcascade_frontalface.xml") ), |
||||
testing::Values( string("cv/shared/lena.png"), |
||||
string("cv/cascadeandhog/images/bttf301.png"), |
||||
string("cv/cascadeandhog/images/class57.png") ), |
||||
testing::Values(30, 64, 90) ) ) |
||||
{ |
||||
const string cascasePath = get<0>(GetParam()); |
||||
const string imagePath = get<1>(GetParam()); |
||||
int min_size = get<2>(GetParam()); |
||||
Size minSize(min_size, min_size); |
||||
|
||||
CascadeClassifier cc( getDataPath(cascasePath) ); |
||||
if (cc.empty()) |
||||
FAIL() << "Can't load cascade file: " << getDataPath(cascasePath); |
||||
|
||||
Mat img = imread(getDataPath(imagePath), IMREAD_GRAYSCALE); |
||||
if (img.empty()) |
||||
FAIL() << "Can't load source image: " << getDataPath(imagePath); |
||||
|
||||
vector<Rect> faces; |
||||
|
||||
equalizeHist(img, img); |
||||
declare.in(img).time(60); |
||||
|
||||
UMat uimg = img.getUMat(ACCESS_READ); |
||||
|
||||
while(next()) |
||||
{ |
||||
faces.clear(); |
||||
cvtest::ocl::perf::safeFinish(); |
||||
|
||||
startTimer(); |
||||
cc.detectMultiScale(uimg, faces, 1.1, 3, 0, minSize); |
||||
stopTimer(); |
||||
} |
||||
|
||||
sort(faces.begin(), faces.end(), comparators::RectLess()); |
||||
SANITY_CHECK(faces, min_size/5); |
||||
} |
||||
|
||||
#endif //HAVE_OPENCL
|
||||
@ -1,56 +0,0 @@ |
||||
#include "perf_precomp.hpp" |
||||
#include <opencv2/imgproc.hpp> |
||||
|
||||
using namespace std; |
||||
using namespace cv; |
||||
using namespace perf; |
||||
using std::tr1::make_tuple; |
||||
using std::tr1::get; |
||||
|
||||
typedef std::tr1::tuple<std::string, int> ImageName_MinSize_t; |
||||
typedef perf::TestBaseWithParam<ImageName_MinSize_t> ImageName_MinSize; |
||||
|
||||
PERF_TEST_P(ImageName_MinSize, CascadeClassifierLBPFrontalFace, |
||||
testing::Combine(testing::Values( std::string("cv/shared/lena.png"), |
||||
std::string("cv/shared/1_itseez-0000289.png"), |
||||
std::string("cv/shared/1_itseez-0000492.png"), |
||||
std::string("cv/shared/1_itseez-0000573.png")), |
||||
testing::Values(24, 30, 40, 50, 60, 70, 80, 90) |
||||
) |
||||
) |
||||
{ |
||||
const string filename = get<0>(GetParam()); |
||||
int min_size = get<1>(GetParam()); |
||||
Size minSize(min_size, min_size); |
||||
|
||||
CascadeClassifier cc(getDataPath("cv/cascadeandhog/cascades/lbpcascade_frontalface.xml")); |
||||
if (cc.empty()) |
||||
FAIL() << "Can't load cascade file"; |
||||
|
||||
Mat img = imread(getDataPath(filename), 0); |
||||
if (img.empty()) |
||||
FAIL() << "Can't load source image"; |
||||
|
||||
vector<Rect> faces; |
||||
|
||||
equalizeHist(img, img); |
||||
declare.in(img); |
||||
|
||||
while(next()) |
||||
{ |
||||
faces.clear(); |
||||
|
||||
startTimer(); |
||||
cc.detectMultiScale(img, faces, 1.1, 3, 0, minSize); |
||||
stopTimer(); |
||||
} |
||||
// for some reason OpenCL version detects the face, which CPU version does not detect, we just remove it
|
||||
// TODO better solution: implement smart way of comparing two set of rectangles
|
||||
if( filename == "cv/shared/1_itseez-0000492.png" && faces.size() == (size_t)3 ) |
||||
{ |
||||
faces.erase(faces.begin()); |
||||
} |
||||
|
||||
std::sort(faces.begin(), faces.end(), comparators::RectLess()); |
||||
SANITY_CHECK(faces, 3.001 * faces.size()); |
||||
} |
||||
Loading…
Reference in new issue