From 089a835c0a008639535b684a9bdd2e0da44f519e Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 26 Mar 2012 18:07:03 +0000 Subject: [PATCH] fixed octave computation in SURF_GPU used random images in gpu filter tests --- modules/gpu/include/opencv2/gpu/gpu.hpp | 15 +- modules/gpu/src/cuda/surf.cu | 13 +- modules/gpu/src/filtering.cpp | 4 +- modules/gpu/src/surf.cpp | 80 ++---- modules/gpu/test/test_core.cpp | 4 +- modules/gpu/test/test_features2d.cpp | 2 +- modules/gpu/test/test_filters.cpp | 467 ++++++++++---------------------- 7 files changed, 200 insertions(+), 385 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 0e6dda9..08c6be5 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1516,13 +1516,14 @@ class CV_EXPORTS SURF_GPU public: enum KeypointLayout { - SF_X = 0, - SF_Y, - SF_LAPLACIAN, - SF_SIZE, - SF_DIR, - SF_HESSIAN, - SF_FEATURE_STRIDE + X_ROW = 0, + Y_ROW, + LAPLACIAN_ROW, + OCTAVE_ROW, + SIZE_ROW, + ANGLE_ROW, + HESSIAN_ROW, + ROWS_COUNT }; //! the default constructor diff --git a/modules/gpu/src/cuda/surf.cu b/modules/gpu/src/cuda/surf.cu index 9d265cc..76181bc 100644 --- a/modules/gpu/src/cuda/surf.cu +++ b/modules/gpu/src/cuda/surf.cu @@ -117,7 +117,7 @@ namespace cv { namespace gpu { namespace device template __device__ float icvCalcHaarPatternSum(const float src[][5], int oldSize, int newSize, int y, int x) { #if __CUDA_ARCH__ >= 200 - typedef double real_t; + typedef double real_t; #else typedef float real_t; #endif @@ -248,7 +248,7 @@ namespace cv { namespace gpu { namespace device template __global__ void icvFindMaximaInLayer(const PtrStepf det, const PtrStepf trace, int4* maxPosBuffer, unsigned int* maxCounter) { - #if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 110 + #if __CUDA_ARCH__ >= 110 extern __shared__ float N9[]; @@ -368,10 +368,10 @@ namespace cv { namespace gpu { namespace device // INTERPOLATION __global__ void icvInterpolateKeypoint(const PtrStepf det, const int4* maxPosBuffer, - float* featureX, float* featureY, int* featureLaplacian, float* featureSize, float* featureHessian, + float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian, unsigned int* featureCounter) { - #if defined (__CUDA_ARCH__) && __CUDA_ARCH__ >= 110 + #if __CUDA_ARCH__ >= 110 const int4 maxPos = maxPosBuffer[blockIdx.x]; @@ -459,6 +459,7 @@ namespace cv { namespace gpu { namespace device featureX[ind] = px; featureY[ind] = py; featureLaplacian[ind] = maxPos.w; + featureOctave[ind] = c_octave; featureSize[ind] = psize; featureHessian[ind] = N9[1][1][1]; } @@ -471,7 +472,7 @@ namespace cv { namespace gpu { namespace device } void icvInterpolateKeypoint_gpu(const PtrStepf& det, const int4* maxPosBuffer, unsigned int maxCounter, - float* featureX, float* featureY, int* featureLaplacian, float* featureSize, float* featureHessian, + float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian, unsigned int* featureCounter) { dim3 threads; @@ -482,7 +483,7 @@ namespace cv { namespace gpu { namespace device dim3 grid; grid.x = maxCounter; - icvInterpolateKeypoint<<>>(det, maxPosBuffer, featureX, featureY, featureLaplacian, featureSize, featureHessian, featureCounter); + icvInterpolateKeypoint<<>>(det, maxPosBuffer, featureX, featureY, featureLaplacian, featureOctave, featureSize, featureHessian, featureCounter); cudaSafeCall( cudaGetLastError() ); cudaSafeCall( cudaDeviceSynchronize() ); diff --git a/modules/gpu/src/filtering.cpp b/modules/gpu/src/filtering.cpp index 4a5f032..ac61d20 100644 --- a/modules/gpu/src/filtering.cpp +++ b/modules/gpu/src/filtering.cpp @@ -948,7 +948,9 @@ namespace { DeviceInfo devInfo; int cc = devInfo.majorVersion() * 10 + devInfo.minorVersion(); - CV_Assert(cc >= 20 || ksize <= 16); + if (ksize > 16 && cc < 20) + CV_Error(CV_StsNotImplemented, "column linear filter doesn't implemented for kernel size > 16 for device with compute capabilities less than 2.0"); + func(src, dst, kernel.ptr(), ksize, anchor, brd_type, cc, StreamAccessor::getStream(s)); } diff --git a/modules/gpu/src/surf.cpp b/modules/gpu/src/surf.cpp index 2bc9f58..dea51a9 100644 --- a/modules/gpu/src/surf.cpp +++ b/modules/gpu/src/surf.cpp @@ -80,7 +80,7 @@ namespace cv { namespace gpu { namespace device int img_rows, int img_cols, int octave, bool use_mask, int nLayers); void icvInterpolateKeypoint_gpu(const PtrStepf& det, const int4* maxPosBuffer, unsigned int maxCounter, - float* featureX, float* featureY, int* featureLaplacian, float* featureSize, float* featureHessian, + float* featureX, float* featureY, int* featureLaplacian, int* featureOctave, float* featureSize, float* featureHessian, unsigned int* featureCounter); void icvCalcOrientation_gpu(const float* featureX, const float* featureY, const float* featureSize, float* featureDir, int nFeatures); @@ -161,7 +161,7 @@ namespace ensureSizeIsEnough(img_rows * (surf_.nOctaveLayers + 2), img_cols, CV_32FC1, surf_.trace); ensureSizeIsEnough(1, maxCandidates, CV_32SC4, surf_.maxPosBuffer); - ensureSizeIsEnough(SURF_GPU::SF_FEATURE_STRIDE, maxFeatures, CV_32FC1, keypoints); + ensureSizeIsEnough(SURF_GPU::ROWS_COUNT, maxFeatures, CV_32FC1, keypoints); keypoints.setTo(Scalar::all(0)); for (int octave = 0; octave < surf_.nOctaves; ++octave) @@ -183,9 +183,10 @@ namespace if (maxCounter > 0) { icvInterpolateKeypoint_gpu(surf_.det, surf_.maxPosBuffer.ptr(), maxCounter, - keypoints.ptr(SURF_GPU::SF_X), keypoints.ptr(SURF_GPU::SF_Y), - keypoints.ptr(SURF_GPU::SF_LAPLACIAN), keypoints.ptr(SURF_GPU::SF_SIZE), - keypoints.ptr(SURF_GPU::SF_HESSIAN), counters.ptr()); + keypoints.ptr(SURF_GPU::X_ROW), keypoints.ptr(SURF_GPU::Y_ROW), + keypoints.ptr(SURF_GPU::LAPLACIAN_ROW), keypoints.ptr(SURF_GPU::OCTAVE_ROW), + keypoints.ptr(SURF_GPU::SIZE_ROW), keypoints.ptr(SURF_GPU::HESSIAN_ROW), + counters.ptr()); } } unsigned int featureCounter; @@ -195,7 +196,7 @@ namespace keypoints.cols = featureCounter; if (surf_.upright) - keypoints.row(SURF_GPU::SF_DIR).setTo(Scalar::all(90.0)); + keypoints.row(SURF_GPU::ANGLE_ROW).setTo(Scalar::all(90.0)); else findOrientation(keypoints); } @@ -205,8 +206,8 @@ namespace const int nFeatures = keypoints.cols; if (nFeatures > 0) { - icvCalcOrientation_gpu(keypoints.ptr(SURF_GPU::SF_X), keypoints.ptr(SURF_GPU::SF_Y), - keypoints.ptr(SURF_GPU::SF_SIZE), keypoints.ptr(SURF_GPU::SF_DIR), nFeatures); + icvCalcOrientation_gpu(keypoints.ptr(SURF_GPU::X_ROW), keypoints.ptr(SURF_GPU::Y_ROW), + keypoints.ptr(SURF_GPU::SIZE_ROW), keypoints.ptr(SURF_GPU::ANGLE_ROW), nFeatures); } } @@ -216,8 +217,8 @@ namespace if (nFeatures > 0) { ensureSizeIsEnough(nFeatures, descriptorSize, CV_32F, descriptors); - compute_descriptors_gpu(descriptors, keypoints.ptr(SURF_GPU::SF_X), keypoints.ptr(SURF_GPU::SF_Y), - keypoints.ptr(SURF_GPU::SF_SIZE), keypoints.ptr(SURF_GPU::SF_DIR), nFeatures); + compute_descriptors_gpu(descriptors, keypoints.ptr(SURF_GPU::X_ROW), keypoints.ptr(SURF_GPU::Y_ROW), + keypoints.ptr(SURF_GPU::SIZE_ROW), keypoints.ptr(SURF_GPU::ANGLE_ROW), nFeatures); } } @@ -266,20 +267,22 @@ void cv::gpu::SURF_GPU::uploadKeypoints(const vector& keypoints, GpuMa keypointsGPU.release(); else { - Mat keypointsCPU(SURF_GPU::SF_FEATURE_STRIDE, static_cast(keypoints.size()), CV_32FC1); + Mat keypointsCPU(SURF_GPU::ROWS_COUNT, static_cast(keypoints.size()), CV_32FC1); - float* kp_x = keypointsCPU.ptr(SURF_GPU::SF_X); - float* kp_y = keypointsCPU.ptr(SURF_GPU::SF_Y); - int* kp_laplacian = keypointsCPU.ptr(SURF_GPU::SF_LAPLACIAN); - float* kp_size = keypointsCPU.ptr(SURF_GPU::SF_SIZE); - float* kp_dir = keypointsCPU.ptr(SURF_GPU::SF_DIR); - float* kp_hessian = keypointsCPU.ptr(SURF_GPU::SF_HESSIAN); + float* kp_x = keypointsCPU.ptr(SURF_GPU::X_ROW); + float* kp_y = keypointsCPU.ptr(SURF_GPU::Y_ROW); + int* kp_laplacian = keypointsCPU.ptr(SURF_GPU::LAPLACIAN_ROW); + int* kp_octave = keypointsCPU.ptr(SURF_GPU::OCTAVE_ROW); + float* kp_size = keypointsCPU.ptr(SURF_GPU::SIZE_ROW); + float* kp_dir = keypointsCPU.ptr(SURF_GPU::ANGLE_ROW); + float* kp_hessian = keypointsCPU.ptr(SURF_GPU::HESSIAN_ROW); for (size_t i = 0, size = keypoints.size(); i < size; ++i) { const KeyPoint& kp = keypoints[i]; kp_x[i] = kp.pt.x; kp_y[i] = kp.pt.y; + kp_octave[i] = kp.octave; kp_size[i] = kp.size; kp_dir[i] = kp.angle; kp_hessian[i] = kp.response; @@ -290,30 +293,6 @@ void cv::gpu::SURF_GPU::uploadKeypoints(const vector& keypoints, GpuMa } } -namespace -{ - int getPointOctave(float size, const SURF_GPU& params) - { - int best_octave = 0; - float min_diff = numeric_limits::max(); - for (int octave = 1; octave < params.nOctaves; ++octave) - { - for (int layer = 0; layer < params.nOctaveLayers; ++layer) - { - float diff = std::abs(size - (float)calcSize(octave, layer)); - if (min_diff > diff) - { - min_diff = diff; - best_octave = octave; - if (min_diff == 0) - return best_octave; - } - } - } - return best_octave; - } -} - void cv::gpu::SURF_GPU::downloadKeypoints(const GpuMat& keypointsGPU, vector& keypoints) { const int nFeatures = keypointsGPU.cols; @@ -322,18 +301,19 @@ void cv::gpu::SURF_GPU::downloadKeypoints(const GpuMat& keypointsGPU, vector(SF_X); - float* kp_y = keypointsCPU.ptr(SF_Y); - int* kp_laplacian = keypointsCPU.ptr(SF_LAPLACIAN); - float* kp_size = keypointsCPU.ptr(SF_SIZE); - float* kp_dir = keypointsCPU.ptr(SF_DIR); - float* kp_hessian = keypointsCPU.ptr(SF_HESSIAN); + + float* kp_x = keypointsCPU.ptr(SURF_GPU::X_ROW); + float* kp_y = keypointsCPU.ptr(SURF_GPU::Y_ROW); + int* kp_laplacian = keypointsCPU.ptr(SURF_GPU::LAPLACIAN_ROW); + int* kp_octave = keypointsCPU.ptr(SURF_GPU::OCTAVE_ROW); + float* kp_size = keypointsCPU.ptr(SURF_GPU::SIZE_ROW); + float* kp_dir = keypointsCPU.ptr(SURF_GPU::ANGLE_ROW); + float* kp_hessian = keypointsCPU.ptr(SURF_GPU::HESSIAN_ROW); for (int i = 0; i < nFeatures; ++i) { @@ -341,10 +321,10 @@ void cv::gpu::SURF_GPU::downloadKeypoints(const GpuMat& keypointsGPU, vector= CV_32F || depth.second >= CV_32F ? 1e-4 : 0.0); + EXPECT_MAT_NEAR(dst_gold, dst, 1.0); } } @@ -2715,7 +2715,7 @@ TEST_P(Sum, Sqr) cv::Scalar val_gold = sqrSumGold(src); - EXPECT_SCALAR_NEAR(val_gold, val, CV_MAT_DEPTH(type) < CV_32F ? 0.0 : 10); + EXPECT_SCALAR_NEAR(val_gold, val, CV_MAT_DEPTH(type) < CV_32F ? 0.0 : 0.5); } INSTANTIATE_TEST_CASE_P(GPU_Core, Sum, testing::Combine( diff --git a/modules/gpu/test/test_features2d.cpp b/modules/gpu/test/test_features2d.cpp index 14642ab..2121820 100644 --- a/modules/gpu/test/test_features2d.cpp +++ b/modules/gpu/test/test_features2d.cpp @@ -85,7 +85,7 @@ testing::AssertionResult assertKeyPointsEquals(const char* gold_expr, const char std::sort(actual.begin(), actual.end(), KeyPointLess()); std::sort(gold.begin(), gold.end(), KeyPointLess()); - for (size_t i; i < gold.size(); ++i) + for (size_t i = 0; i < gold.size(); ++i) { const cv::KeyPoint& p1 = gold[i]; const cv::KeyPoint& p2 = actual[i]; diff --git a/modules/gpu/test/test_filters.cpp b/modules/gpu/test/test_filters.cpp index 1424165..6b9f23c 100644 --- a/modules/gpu/test/test_filters.cpp +++ b/modules/gpu/test/test_filters.cpp @@ -68,59 +68,45 @@ cv::Mat getInnerROI(cv::InputArray m, int ksize) IMPLEMENT_PARAM_CLASS(Anchor, cv::Point) -PARAM_TEST_CASE(Blur, cv::gpu::DeviceInfo, KSize, Anchor, UseRoi) +PARAM_TEST_CASE(Blur, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, Anchor, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; cv::Size ksize; cv::Point anchor; bool useRoi; - cv::Mat img; - virtual void SetUp() { devInfo = GET_PARAM(0); - ksize = GET_PARAM(1); - anchor = GET_PARAM(2); - useRoi = GET_PARAM(3); + size = GET_PARAM(1); + type = GET_PARAM(2); + ksize = GET_PARAM(3); + anchor = GET_PARAM(4); + useRoi = GET_PARAM(5); cv::gpu::setDevice(devInfo.deviceID()); - - img = readImage("stereobp/aloe-L.png"); - ASSERT_FALSE(img.empty()); } }; -TEST_P(Blur, Gray) -{ - cv::Mat img_gray; - cv::cvtColor(img, img_gray, CV_BGR2GRAY); - - cv::gpu::GpuMat dst; - cv::gpu::blur(loadMat(img_gray, useRoi), dst, ksize, anchor); - - cv::Mat dst_gold; - cv::blur(img_gray, dst_gold, ksize, anchor); - - EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0); -} - -TEST_P(Blur, Color) +TEST_P(Blur, Accuracy) { - cv::Mat img_rgba; - cv::cvtColor(img, img_rgba, CV_BGR2BGRA); + cv::Mat src = randomMat(size, type); - cv::gpu::GpuMat dst; - cv::gpu::blur(loadMat(img_rgba, useRoi), dst, ksize, anchor); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); + cv::gpu::blur(loadMat(src, useRoi), dst, ksize, anchor); cv::Mat dst_gold; - cv::blur(img_rgba, dst_gold, ksize, anchor); + cv::blur(src, dst_gold, ksize, anchor); EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 1.0); } INSTANTIATE_TEST_CASE_P(GPU_Filter, Blur, testing::Combine( ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), testing::Values(KSize(3, 3), KSize(5, 5), KSize(7, 7)), testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))), WHOLE_SUBMAT)); @@ -131,69 +117,52 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Blur, testing::Combine( IMPLEMENT_PARAM_CLASS(Deriv_X, int) IMPLEMENT_PARAM_CLASS(Deriv_Y, int) -PARAM_TEST_CASE(Sobel, cv::gpu::DeviceInfo, KSize, Deriv_X, Deriv_Y, BorderType, UseRoi) +PARAM_TEST_CASE(Sobel, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, Deriv_X, Deriv_Y, BorderType, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; cv::Size ksize; int dx; int dy; int borderType; bool useRoi; - cv::Mat img; - virtual void SetUp() { devInfo = GET_PARAM(0); - ksize = GET_PARAM(1); - dx = GET_PARAM(2); - dy = GET_PARAM(3); - borderType = GET_PARAM(4); - useRoi = GET_PARAM(5); + size = GET_PARAM(1); + type = GET_PARAM(2); + ksize = GET_PARAM(3); + dx = GET_PARAM(4); + dy = GET_PARAM(5); + borderType = GET_PARAM(6); + useRoi = GET_PARAM(7); cv::gpu::setDevice(devInfo.deviceID()); - - img = readImage("stereobp/aloe-L.png"); - ASSERT_FALSE(img.empty()); } }; -TEST_P(Sobel, Gray) -{ - if (dx == 0 && dy == 0) - return; - - cv::Mat img_gray; - cv::cvtColor(img, img_gray, CV_BGR2GRAY); - - cv::gpu::GpuMat dst; - cv::gpu::Sobel(loadMat(img_gray, useRoi), dst, -1, dx, dy, ksize.width, 1.0, borderType); - - cv::Mat dst_gold; - cv::Sobel(img_gray, dst_gold, -1, dx, dy, ksize.width, 1.0, 0.0, borderType); - - EXPECT_MAT_NEAR(dst_gold, dst, 0.0); -} - -TEST_P(Sobel, Color) +TEST_P(Sobel, Accuracy) { if (dx == 0 && dy == 0) return; - cv::Mat img_rgba; - cv::cvtColor(img, img_rgba, CV_BGR2BGRA); + cv::Mat src = randomMat(size, type); - cv::gpu::GpuMat dst; - cv::gpu::Sobel(loadMat(img_rgba, useRoi), dst, -1, dx, dy, ksize.width, 1.0, borderType); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); + cv::gpu::Sobel(loadMat(src, useRoi), dst, -1, dx, dy, ksize.width, 1.0, borderType); cv::Mat dst_gold; - cv::Sobel(img_rgba, dst_gold, -1, dx, dy, ksize.width, 1.0, 0.0, borderType); + cv::Sobel(src, dst_gold, -1, dx, dy, ksize.width, 1.0, 0.0, borderType); EXPECT_MAT_NEAR(dst_gold, dst, 0.0); } INSTANTIATE_TEST_CASE_P(GPU_Filter, Sobel, testing::Combine( ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), testing::Values(KSize(3, 3), KSize(5, 5), KSize(7, 7)), testing::Values(Deriv_X(0), Deriv_X(1), Deriv_X(2)), testing::Values(Deriv_Y(0), Deriv_Y(1), Deriv_Y(2)), @@ -206,67 +175,50 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Sobel, testing::Combine( ///////////////////////////////////////////////////////////////////////////////////////////////// // Scharr -PARAM_TEST_CASE(Scharr, cv::gpu::DeviceInfo, Deriv_X, Deriv_Y, BorderType, UseRoi) +PARAM_TEST_CASE(Scharr, cv::gpu::DeviceInfo, cv::Size, MatType, Deriv_X, Deriv_Y, BorderType, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; int dx; int dy; int borderType; bool useRoi; - cv::Mat img; - virtual void SetUp() { devInfo = GET_PARAM(0); - dx = GET_PARAM(1); - dy = GET_PARAM(2); - borderType = GET_PARAM(3); - useRoi = GET_PARAM(4); + size = GET_PARAM(1); + type = GET_PARAM(2); + dx = GET_PARAM(3); + dy = GET_PARAM(4); + borderType = GET_PARAM(5); + useRoi = GET_PARAM(6); cv::gpu::setDevice(devInfo.deviceID()); - - img = readImage("stereobp/aloe-L.png"); - ASSERT_FALSE(img.empty()); } }; -TEST_P(Scharr, Gray) -{ - if (dx + dy != 1) - return; - - cv::Mat img_gray; - cv::cvtColor(img, img_gray, CV_BGR2GRAY); - - cv::gpu::GpuMat dst; - cv::gpu::Scharr(loadMat(img_gray, useRoi), dst, -1, dx, dy, 1.0, borderType); - - cv::Mat dst_gold; - cv::Scharr(img_gray, dst_gold, -1, dx, dy, 1.0, 0.0, borderType); - - EXPECT_MAT_NEAR(dst_gold, dst, 0.0); -} - -TEST_P(Scharr, Color) +TEST_P(Scharr, Accuracy) { if (dx + dy != 1) return; - cv::Mat img_rgba; - cv::cvtColor(img, img_rgba, CV_BGR2BGRA); + cv::Mat src = randomMat(size, type); - cv::gpu::GpuMat dst; - cv::gpu::Scharr(loadMat(img_rgba, useRoi), dst, -1, dx, dy, 1.0, borderType); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); + cv::gpu::Scharr(loadMat(src, useRoi), dst, -1, dx, dy, 1.0, borderType); cv::Mat dst_gold; - cv::Scharr(img_rgba, dst_gold, -1, dx, dy, 1.0, 0.0, borderType); + cv::Scharr(src, dst_gold, -1, dx, dy, 1.0, 0.0, borderType); EXPECT_MAT_NEAR(dst_gold, dst, 0.0); } INSTANTIATE_TEST_CASE_P(GPU_Filter, Scharr, testing::Combine( ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), testing::Values(Deriv_X(0), Deriv_X(1)), testing::Values(Deriv_Y(0), Deriv_Y(1)), testing::Values(BorderType(cv::BORDER_REFLECT101), @@ -278,63 +230,47 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Scharr, testing::Combine( ///////////////////////////////////////////////////////////////////////////////////////////////// // GaussianBlur -PARAM_TEST_CASE(GaussianBlur, cv::gpu::DeviceInfo, KSize, BorderType, UseRoi) +PARAM_TEST_CASE(GaussianBlur, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, BorderType, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; cv::Size ksize; int borderType; bool useRoi; - cv::Mat img; - double sigma1, sigma2; - virtual void SetUp() { devInfo = GET_PARAM(0); - ksize = GET_PARAM(1); - borderType = GET_PARAM(2); - useRoi = GET_PARAM(3); + size = GET_PARAM(1); + type = GET_PARAM(2); + ksize = GET_PARAM(3); + borderType = GET_PARAM(4); + useRoi = GET_PARAM(5); cv::gpu::setDevice(devInfo.deviceID()); - - img = readImage("stereobp/aloe-L.png"); - ASSERT_FALSE(img.empty()); - - sigma1 = randomDouble(0.1, 1.0); - sigma2 = randomDouble(0.1, 1.0); } }; -TEST_P(GaussianBlur, Gray) -{ - cv::Mat img_gray; - cv::cvtColor(img, img_gray, CV_BGR2GRAY); - - cv::gpu::GpuMat dst; - cv::gpu::GaussianBlur(loadMat(img_gray, useRoi), dst, ksize, sigma1, sigma2, borderType); - - cv::Mat dst_gold; - cv::GaussianBlur(img_gray, dst_gold, ksize, sigma1, sigma2, borderType); - - EXPECT_MAT_NEAR(dst_gold, dst, 4.0); -} - -TEST_P(GaussianBlur, Color) +TEST_P(GaussianBlur, Accuracy) { - cv::Mat img_rgba; - cv::cvtColor(img, img_rgba, CV_BGR2BGRA); + cv::Mat src = randomMat(size, type); + double sigma1 = randomDouble(0.1, 1.0); + double sigma2 = randomDouble(0.1, 1.0); - cv::gpu::GpuMat dst; - cv::gpu::GaussianBlur(loadMat(img_rgba, useRoi), dst, ksize, sigma1, sigma2, borderType); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); + cv::gpu::GaussianBlur(loadMat(src, useRoi), dst, ksize, sigma1, sigma2, borderType); cv::Mat dst_gold; - cv::GaussianBlur(img_rgba, dst_gold, ksize, sigma1, sigma2, borderType); + cv::GaussianBlur(src, dst_gold, ksize, sigma1, sigma2, borderType); EXPECT_MAT_NEAR(dst_gold, dst, 4.0); } INSTANTIATE_TEST_CASE_P(GPU_Filter, GaussianBlur, testing::Combine( ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), testing::Values(KSize(3, 3), KSize(5, 5), KSize(7, 7), @@ -359,72 +295,46 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, GaussianBlur, testing::Combine( ///////////////////////////////////////////////////////////////////////////////////////////////// // Laplacian -PARAM_TEST_CASE(Laplacian, cv::gpu::DeviceInfo, KSize, UseRoi) +PARAM_TEST_CASE(Laplacian, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; cv::Size ksize; bool useRoi; - cv::Mat img; - virtual void SetUp() { devInfo = GET_PARAM(0); - ksize = GET_PARAM(1); - useRoi = GET_PARAM(2); + size = GET_PARAM(1); + type = GET_PARAM(2); + ksize = GET_PARAM(3); + useRoi = GET_PARAM(4); cv::gpu::setDevice(devInfo.deviceID()); - - img = readImage("stereobp/aloe-L.png"); - ASSERT_FALSE(img.empty()); } }; -TEST_P(Laplacian, Gray) -{ - cv::Mat img_gray; - cv::cvtColor(img, img_gray, CV_BGR2GRAY); - - cv::gpu::GpuMat dst; - cv::gpu::Laplacian(loadMat(img_gray, useRoi), dst, -1, ksize.width); - - cv::Mat dst_gold; - cv::Laplacian(img_gray, dst_gold, -1, ksize.width); - - EXPECT_MAT_NEAR(getInnerROI(dst_gold, cv::Size(3, 3)), getInnerROI(dst, cv::Size(3, 3)), 0.0); -} - -TEST_P(Laplacian, Color) +TEST_P(Laplacian, Accuracy) { - cv::Mat img_rgba; - cv::cvtColor(img, img_rgba, CV_BGR2BGRA); + cv::Mat src = randomMat(size, type); - cv::gpu::GpuMat dst; - cv::gpu::Laplacian(loadMat(img_rgba, useRoi), dst, -1, ksize.width); - - cv::Mat dst_gold; - cv::Laplacian(img_rgba, dst_gold, -1, ksize.width); - - EXPECT_MAT_NEAR(getInnerROI(dst_gold, cv::Size(3, 3)), getInnerROI(dst, cv::Size(3, 3)), 0.0); -} - -TEST_P(Laplacian, Gray_32FC1) -{ - cv::Mat src; - cv::cvtColor(img, src, CV_BGR2GRAY); - src.convertTo(src, CV_32F, 1.0 / 255.0); - - cv::gpu::GpuMat dst; + cv::gpu::GpuMat dst = createMat(size, type, useRoi); cv::gpu::Laplacian(loadMat(src, useRoi), dst, -1, ksize.width); cv::Mat dst_gold; cv::Laplacian(src, dst_gold, -1, ksize.width); - EXPECT_MAT_NEAR(dst_gold, dst, 0.0); + if (type == CV_32FC1) + EXPECT_MAT_NEAR(dst_gold, dst, 0.0); + else + EXPECT_MAT_NEAR(getInnerROI(dst_gold, cv::Size(3, 3)), getInnerROI(dst, cv::Size(3, 3)), 0.0); } INSTANTIATE_TEST_CASE_P(GPU_Filter, Laplacian, testing::Combine( ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)), testing::Values(KSize(1, 1), KSize(3, 3)), WHOLE_SUBMAT)); @@ -433,58 +343,38 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Laplacian, testing::Combine( IMPLEMENT_PARAM_CLASS(Iterations, int) -PARAM_TEST_CASE(Erode, cv::gpu::DeviceInfo, Anchor, Iterations, UseRoi) +PARAM_TEST_CASE(Erode, cv::gpu::DeviceInfo, cv::Size, MatType, Anchor, Iterations, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; cv::Point anchor; int iterations; bool useRoi; - cv::Mat img; - cv::Mat kernel; - virtual void SetUp() { devInfo = GET_PARAM(0); - anchor = GET_PARAM(1); - iterations = GET_PARAM(2); - useRoi = GET_PARAM(3); + size = GET_PARAM(1); + type = GET_PARAM(2); + anchor = GET_PARAM(3); + iterations = GET_PARAM(4); + useRoi = GET_PARAM(5); cv::gpu::setDevice(devInfo.deviceID()); - - img = readImage("stereobp/aloe-L.png"); - ASSERT_FALSE(img.empty()); - - kernel = cv::Mat::ones(3, 3, CV_8U); } }; -TEST_P(Erode, Gray) -{ - cv::Mat img_gray; - cv::cvtColor(img, img_gray, CV_BGR2GRAY); - - cv::gpu::GpuMat dst; - cv::gpu::erode(loadMat(img_gray, useRoi), dst, kernel, anchor, iterations); - - cv::Mat dst_gold; - cv::erode(img_gray, dst_gold, kernel, anchor, iterations); - - cv::Size ksize = cv::Size(kernel.cols + iterations * (kernel.cols - 1), kernel.rows + iterations * (kernel.rows - 1)); - - EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0); -} - -TEST_P(Erode, Color) +TEST_P(Erode, Accuracy) { - cv::Mat img_rgba; - cv::cvtColor(img, img_rgba, CV_BGR2BGRA); + cv::Mat src = randomMat(size, type); + cv::Mat kernel = cv::Mat::ones(3, 3, CV_8U); - cv::gpu::GpuMat dst; - cv::gpu::erode(loadMat(img_rgba, useRoi), dst, kernel, anchor, iterations); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); + cv::gpu::erode(loadMat(src, useRoi), dst, kernel, anchor, iterations); cv::Mat dst_gold; - cv::erode(img_rgba, dst_gold, kernel, anchor, iterations); + cv::erode(src, dst_gold, kernel, anchor, iterations); cv::Size ksize = cv::Size(kernel.cols + iterations * (kernel.cols - 1), kernel.rows + iterations * (kernel.rows - 1)); @@ -493,6 +383,8 @@ TEST_P(Erode, Color) INSTANTIATE_TEST_CASE_P(GPU_Filter, Erode, testing::Combine( ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))), testing::Values(Iterations(1), Iterations(2), Iterations(3)), WHOLE_SUBMAT)); @@ -500,58 +392,38 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Erode, testing::Combine( ///////////////////////////////////////////////////////////////////////////////////////////////// // Dilate -PARAM_TEST_CASE(Dilate, cv::gpu::DeviceInfo, Anchor, Iterations, UseRoi) +PARAM_TEST_CASE(Dilate, cv::gpu::DeviceInfo, cv::Size, MatType, Anchor, Iterations, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; cv::Point anchor; int iterations; bool useRoi; - cv::Mat img; - cv::Mat kernel; - virtual void SetUp() { devInfo = GET_PARAM(0); - anchor = GET_PARAM(1); - iterations = GET_PARAM(2); - useRoi = GET_PARAM(3); + size = GET_PARAM(1); + type = GET_PARAM(2); + anchor = GET_PARAM(3); + iterations = GET_PARAM(4); + useRoi = GET_PARAM(5); cv::gpu::setDevice(devInfo.deviceID()); - - img = readImage("stereobp/aloe-L.png"); - ASSERT_FALSE(img.empty()); - - kernel = cv::Mat::ones(3, 3, CV_8U); } }; -TEST_P(Dilate, Gray) -{ - cv::Mat img_gray; - cv::cvtColor(img, img_gray, CV_BGR2GRAY); - - cv::gpu::GpuMat dst; - cv::gpu::dilate(loadMat(img_gray, useRoi), dst, kernel, anchor, iterations); - - cv::Mat dst_gold; - cv::dilate(img_gray, dst_gold, kernel, anchor, iterations); - - cv::Size ksize = cv::Size(kernel.cols + iterations * (kernel.cols - 1), kernel.rows + iterations * (kernel.rows - 1)); - - EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0); -} - -TEST_P(Dilate, Color) +TEST_P(Dilate, Accuracy) { - cv::Mat img_rgba; - cv::cvtColor(img, img_rgba, CV_BGR2BGRA); + cv::Mat src = randomMat(size, type); + cv::Mat kernel = cv::Mat::ones(3, 3, CV_8U); - cv::gpu::GpuMat dst; - cv::gpu::dilate(loadMat(img_rgba, useRoi), dst, kernel, anchor, iterations); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); + cv::gpu::dilate(loadMat(src, useRoi), dst, kernel, anchor, iterations); cv::Mat dst_gold; - cv::dilate(img_rgba, dst_gold, kernel, anchor, iterations); + cv::dilate(src, dst_gold, kernel, anchor, iterations); cv::Size ksize = cv::Size(kernel.cols + iterations * (kernel.cols - 1), kernel.rows + iterations * (kernel.rows - 1)); @@ -560,6 +432,8 @@ TEST_P(Dilate, Color) INSTANTIATE_TEST_CASE_P(GPU_Filter, Dilate, testing::Combine( ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))), testing::Values(Iterations(1), Iterations(2), Iterations(3)), WHOLE_SUBMAT)); @@ -570,60 +444,40 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, Dilate, testing::Combine( CV_ENUM(MorphOp, cv::MORPH_OPEN, cv::MORPH_CLOSE, cv::MORPH_GRADIENT, cv::MORPH_TOPHAT, cv::MORPH_BLACKHAT) #define ALL_MORPH_OPS testing::Values(MorphOp(cv::MORPH_OPEN), MorphOp(cv::MORPH_CLOSE), MorphOp(cv::MORPH_GRADIENT), MorphOp(cv::MORPH_TOPHAT), MorphOp(cv::MORPH_BLACKHAT)) -PARAM_TEST_CASE(MorphEx, cv::gpu::DeviceInfo, MorphOp, Anchor, Iterations, UseRoi) +PARAM_TEST_CASE(MorphEx, cv::gpu::DeviceInfo, cv::Size, MatType, MorphOp, Anchor, Iterations, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; int morphOp; cv::Point anchor; int iterations; bool useRoi; - cv::Mat img; - cv::Mat kernel; - virtual void SetUp() { devInfo = GET_PARAM(0); - morphOp = GET_PARAM(1); - anchor = GET_PARAM(2); - iterations = GET_PARAM(3); - useRoi = GET_PARAM(4); + size = GET_PARAM(1); + type = GET_PARAM(2); + morphOp = GET_PARAM(3); + anchor = GET_PARAM(4); + iterations = GET_PARAM(5); + useRoi = GET_PARAM(6); cv::gpu::setDevice(devInfo.deviceID()); - - img = readImage("stereobp/aloe-L.png"); - ASSERT_FALSE(img.empty()); - - kernel = cv::Mat::ones(3, 3, CV_8U); } }; -TEST_P(MorphEx, Gray) -{ - cv::Mat img_gray; - cv::cvtColor(img, img_gray, CV_BGR2GRAY); - - cv::gpu::GpuMat dst; - cv::gpu::morphologyEx(loadMat(img_gray, useRoi), dst, morphOp, kernel, anchor, iterations); - - cv::Mat dst_gold; - cv::morphologyEx(img_gray, dst_gold, morphOp, kernel, anchor, iterations); - - cv::Size border = cv::Size(kernel.cols + (iterations + 1) * kernel.cols + 2, kernel.rows + (iterations + 1) * kernel.rows + 2); - - EXPECT_MAT_NEAR(getInnerROI(dst_gold, border), getInnerROI(dst, border), 0.0); -} - -TEST_P(MorphEx, Color) +TEST_P(MorphEx, Accuracy) { - cv::Mat img_rgba; - cv::cvtColor(img, img_rgba, CV_BGR2BGRA); + cv::Mat src = randomMat(size, type); + cv::Mat kernel = cv::Mat::ones(3, 3, CV_8U); - cv::gpu::GpuMat dst; - cv::gpu::morphologyEx(loadMat(img_rgba, useRoi), dst, morphOp, kernel, anchor, iterations); + cv::gpu::GpuMat dst = createMat(size, type, useRoi); + cv::gpu::morphologyEx(loadMat(src, useRoi), dst, morphOp, kernel, anchor, iterations); cv::Mat dst_gold; - cv::morphologyEx(img_rgba, dst_gold, morphOp, kernel, anchor, iterations); + cv::morphologyEx(src, dst_gold, morphOp, kernel, anchor, iterations); cv::Size border = cv::Size(kernel.cols + (iterations + 1) * kernel.cols + 2, kernel.rows + (iterations + 1) * kernel.rows + 2); @@ -632,6 +486,8 @@ TEST_P(MorphEx, Color) INSTANTIATE_TEST_CASE_P(GPU_Filter, MorphEx, testing::Combine( ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), MatType(CV_8UC4)), ALL_MORPH_OPS, testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))), testing::Values(Iterations(1), Iterations(2), Iterations(3)), @@ -640,9 +496,11 @@ INSTANTIATE_TEST_CASE_P(GPU_Filter, MorphEx, testing::Combine( ///////////////////////////////////////////////////////////////////////////////////////////////// // Filter2D -PARAM_TEST_CASE(Filter2D, cv::gpu::DeviceInfo, KSize, Anchor, UseRoi) +PARAM_TEST_CASE(Filter2D, cv::gpu::DeviceInfo, cv::Size, MatType, KSize, Anchor, UseRoi) { cv::gpu::DeviceInfo devInfo; + cv::Size size; + int type; cv::Size ksize; cv::Point anchor; bool useRoi; @@ -653,64 +511,37 @@ PARAM_TEST_CASE(Filter2D, cv::gpu::DeviceInfo, KSize, Anchor, UseRoi) virtual void SetUp() { devInfo = GET_PARAM(0); - ksize = GET_PARAM(1); - anchor = GET_PARAM(2); - useRoi = GET_PARAM(3); + size = GET_PARAM(1); + type = GET_PARAM(2); + ksize = GET_PARAM(3); + anchor = GET_PARAM(4); + useRoi = GET_PARAM(5); cv::gpu::setDevice(devInfo.deviceID()); - - img = readImage("stereobp/aloe-L.png"); - ASSERT_FALSE(img.empty()); - - kernel = cv::Mat::ones(ksize.height, ksize.width, CV_32FC1); } }; -TEST_P(Filter2D, Gray) -{ - cv::Mat img_gray; - cv::cvtColor(img, img_gray, CV_BGR2GRAY); - - cv::gpu::GpuMat dst; - cv::gpu::filter2D(loadMat(img_gray, useRoi), dst, -1, kernel, anchor); - - cv::Mat dst_gold; - cv::filter2D(img_gray, dst_gold, -1, kernel, anchor, 0, cv::BORDER_CONSTANT); - - EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0); -} - -TEST_P(Filter2D, Color) -{ - cv::Mat img_rgba; - cv::cvtColor(img, img_rgba, CV_BGR2BGRA); - - cv::gpu::GpuMat dst; - cv::gpu::filter2D(loadMat(img_rgba, useRoi), dst, -1, kernel, anchor); - - cv::Mat dst_gold; - cv::filter2D(img_rgba, dst_gold, -1, kernel, anchor, 0, cv::BORDER_CONSTANT); - - EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0); -} - -TEST_P(Filter2D, Gray_32FC1) +TEST_P(Filter2D, Accuracy) { - cv::Mat src; - cv::cvtColor(img, src, CV_BGR2GRAY); - src.convertTo(src, CV_32F, 1.0 / 255.0); + cv::Mat src = randomMat(size, type); + cv::Mat kernel = cv::Mat::ones(ksize.height, ksize.width, CV_32FC1); - cv::gpu::GpuMat dst; + cv::gpu::GpuMat dst = createMat(size, type, useRoi); cv::gpu::filter2D(loadMat(src, useRoi), dst, -1, kernel, anchor); cv::Mat dst_gold; - cv::filter2D(src, dst_gold, -1, kernel, anchor); + cv::filter2D(src, dst_gold, -1, kernel, anchor, 0, type == CV_32FC1 ? cv::BORDER_DEFAULT : cv::BORDER_CONSTANT); - EXPECT_MAT_NEAR(dst_gold, dst, 1e-3); + if (type == CV_32FC1) + EXPECT_MAT_NEAR(dst_gold, dst, 1e-1); + else + EXPECT_MAT_NEAR(getInnerROI(dst_gold, ksize), getInnerROI(dst, ksize), 0.0); } INSTANTIATE_TEST_CASE_P(GPU_Filter, Filter2D, testing::Combine( ALL_DEVICES, + DIFFERENT_SIZES, + testing::Values(MatType(CV_8UC1), MatType(CV_8UC4), MatType(CV_32FC1)), testing::Values(KSize(3, 3), KSize(5, 5), KSize(7, 7), KSize(11, 11), KSize(13, 13), KSize(15, 15)), testing::Values(Anchor(cv::Point(-1, -1)), Anchor(cv::Point(0, 0)), Anchor(cv::Point(2, 2))), WHOLE_SUBMAT)); -- 2.7.4