mirror/opencv
1
0
Fork 0
mirror of https://github.com/opencv/opencv.git synced 2025-01-19 06:53:50 +08:00
Code Issues Actions Packages Projects Releases Wiki Activity

refactored gpu module tests

Browse Source
This commit is contained in:
Vladislav Vinogradov 2011-06-29 10:14:16 +00:00
parent 6e3142f04b
commit d1b4b5f0de
37 changed files with 7383 additions and 5790 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
modules/gpu/include/opencv2/gpu/gpu.hpp
View File

@ -124,6 +124,8 @@ namespace cv
// Checks whether the GPU module can be run on the given device
bool isCompatible() const;
int deviceID() const { return device_id_; }
private:
void query();
void queryMemory(size_t& free_memory, size_t& total_memory) const;
@ -517,14 +519,14 @@ namespace cv
//! supports CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1 types
CV_EXPORTS void multiply(const GpuMat& a, const GpuMat& b, GpuMat& c, Stream& stream = Stream::Null());
//! multiplies matrix to a scalar (c = a * s)
//! supports CV_32FC1 and CV_32FC2 type
//! supports CV_32FC1 type
CV_EXPORTS void multiply(const GpuMat& a, const Scalar& sc, GpuMat& c, Stream& stream = Stream::Null());
//! computes element-wise quotient of the two arrays (c = a / b)
//! supports CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1 types
CV_EXPORTS void divide(const GpuMat& a, const GpuMat& b, GpuMat& c, Stream& stream = Stream::Null());
//! computes element-wise quotient of matrix and scalar (c = a / s)
//! supports CV_32FC1 and CV_32FC2 type
//! supports CV_32FC1 type
CV_EXPORTS void divide(const GpuMat& a, const Scalar& sc, GpuMat& c, Stream& stream = Stream::Null());
//! computes exponent of each matrix element (b = e**a)
@ -1412,9 +1414,9 @@ namespace cv
void radiusMatch(const GpuMat& queryDescs, std::vector< std::vector<DMatch> >& matches, float maxDistance,
const std::vector<GpuMat>& masks = std::vector<GpuMat>(), bool compactResult = false);
private:
DistType distType;
private:
std::vector<GpuMat> trainDescCollection;
};

574
modules/gpu/src/calib3d.cpp
View File

@ -1,290 +1,290 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#if !defined(HAVE_CUDA)
void cv::gpu::transformPoints(const GpuMat&, const Mat&, const Mat&, GpuMat&, Stream&) { throw_nogpu(); }
void cv::gpu::projectPoints(const GpuMat&, const Mat&, const Mat&, const Mat&, const Mat&, GpuMat&, Stream&) { throw_nogpu(); }
void cv::gpu::solvePnPRansac(const Mat&, const Mat&, const Mat&, const Mat&, Mat&, Mat&, bool, int, float, int, vector<int>*) { throw_nogpu(); }
#else
using namespace cv;
using namespace cv::gpu;
namespace cv { namespace gpu { namespace transform_points
{
void call(const DevMem2D_<float3> src, const float* rot, const float* transl, DevMem2D_<float3> dst, cudaStream_t stream);
}}}
namespace
{
void transformPointsCaller(const GpuMat& src, const Mat& rvec, const Mat& tvec, GpuMat& dst, cudaStream_t stream)
{
CV_Assert(src.rows == 1 && src.cols > 0 && src.type() == CV_32FC3);
CV_Assert(rvec.size() == Size(3, 1) && rvec.type() == CV_32F);
CV_Assert(tvec.size() == Size(3, 1) && tvec.type() == CV_32F);
// Convert rotation vector into matrix
Mat rot;
Rodrigues(rvec, rot);
dst.create(src.size(), src.type());
transform_points::call(src, rot.ptr<float>(), tvec.ptr<float>(), dst, stream);
}
}
void cv::gpu::transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec, GpuMat& dst, Stream& stream)
{
::transformPointsCaller(src, rvec, tvec, dst, StreamAccessor::getStream(stream));
}
namespace cv { namespace gpu { namespace project_points
{
void call(const DevMem2D_<float3> src, const float* rot, const float* transl, const float* proj, DevMem2D_<float2> dst, cudaStream_t stream);
}}}
namespace
{
void projectPointsCaller(const GpuMat& src, const Mat& rvec, const Mat& tvec, const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst, cudaStream_t stream)
{
CV_Assert(src.rows == 1 && src.cols > 0 && src.type() == CV_32FC3);
CV_Assert(rvec.size() == Size(3, 1) && rvec.type() == CV_32F);
CV_Assert(tvec.size() == Size(3, 1) && tvec.type() == CV_32F);
CV_Assert(camera_mat.size() == Size(3, 3) && camera_mat.type() == CV_32F);
CV_Assert(dist_coef.empty()); // Undistortion isn't supported
// Convert rotation vector into matrix
Mat rot;
Rodrigues(rvec, rot);
dst.create(src.size(), CV_32FC2);
project_points::call(src, rot.ptr<float>(), tvec.ptr<float>(), camera_mat.ptr<float>(), dst,stream);
}
}
void cv::gpu::projectPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec, const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst, Stream& stream)
{
::projectPointsCaller(src, rvec, tvec, camera_mat, dist_coef, dst, StreamAccessor::getStream(stream));
}
namespace cv { namespace gpu { namespace solve_pnp_ransac
{
int maxNumIters();
void computeHypothesisScores(
const int num_hypotheses, const int num_points, const float* rot_matrices,
const float3* transl_vectors, const float3* object, const float2* image,
const float dist_threshold, int* hypothesis_scores);
}}}
namespace
{
// Selects subset_size random different points from [0, num_points - 1] range
void selectRandom(int subset_size, int num_points, vector<int>& subset)
{
subset.resize(subset_size);
for (int i = 0; i < subset_size; ++i)
{
bool was;
do
{
subset[i] = rand() % num_points;
was = false;
for (int j = 0; j < i; ++j)
if (subset[j] == subset[i])
{
was = true;
break;
}
} while (was);
}
}
// Computes rotation, translation pair for small subsets if the input data
class TransformHypothesesGenerator
{
public:
TransformHypothesesGenerator(const Mat& object_, const Mat& image_, const Mat& dist_coef_,
const Mat& camera_mat_, int num_points_, int subset_size_,
Mat rot_matrices_, Mat transl_vectors_)
: object(&object_), image(&image_), dist_coef(&dist_coef_), camera_mat(&camera_mat_),
num_points(num_points_), subset_size(subset_size_), rot_matrices(rot_matrices_),
transl_vectors(transl_vectors_) {}
void operator()(const BlockedRange& range) const
{
// Input data for generation of the current hypothesis
vector<int> subset_indices(subset_size);
Mat_<Point3f> object_subset(1, subset_size);
Mat_<Point2f> image_subset(1, subset_size);
// Current hypothesis data
Mat rot_vec(1, 3, CV_64F);
Mat rot_mat(3, 3, CV_64F);
Mat transl_vec(1, 3, CV_64F);
for (int iter = range.begin(); iter < range.end(); ++iter)
{
selectRandom(subset_size, num_points, subset_indices);
for (int i = 0; i < subset_size; ++i)
{
object_subset(0, i) = object->at<Point3f>(subset_indices[i]);
image_subset(0, i) = image->at<Point2f>(subset_indices[i]);
}
solvePnP(object_subset, image_subset, *camera_mat, *dist_coef, rot_vec, transl_vec);
// Remember translation vector
Mat transl_vec_ = transl_vectors.colRange(iter * 3, (iter + 1) * 3);
transl_vec = transl_vec.reshape(0, 1);
transl_vec.convertTo(transl_vec_, CV_32F);
// Remember rotation matrix
Rodrigues(rot_vec, rot_mat);
Mat rot_mat_ = rot_matrices.colRange(iter * 9, (iter + 1) * 9).reshape(0, 3);
rot_mat.convertTo(rot_mat_, CV_32F);
}
}
const Mat* object;
const Mat* image;
const Mat* dist_coef;
const Mat* camera_mat;
int num_points;
int subset_size;
// Hypotheses storage (global)
Mat rot_matrices;
Mat transl_vectors;
};
}
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of the copyright holders may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "precomp.hpp"
#if !defined(HAVE_CUDA)
void cv::gpu::transformPoints(const GpuMat&, const Mat&, const Mat&, GpuMat&, Stream&) { throw_nogpu(); }
void cv::gpu::projectPoints(const GpuMat&, const Mat&, const Mat&, const Mat&, const Mat&, GpuMat&, Stream&) { throw_nogpu(); }
void cv::gpu::solvePnPRansac(const Mat&, const Mat&, const Mat&, const Mat&, Mat&, Mat&, bool, int, float, int, vector<int>*) { throw_nogpu(); }
#else
using namespace cv;
using namespace cv::gpu;
namespace cv { namespace gpu { namespace transform_points
{
void call(const DevMem2D_<float3> src, const float* rot, const float* transl, DevMem2D_<float3> dst, cudaStream_t stream);
}}}
namespace
{
void transformPointsCaller(const GpuMat& src, const Mat& rvec, const Mat& tvec, GpuMat& dst, cudaStream_t stream)
{
CV_Assert(src.rows == 1 && src.cols > 0 && src.type() == CV_32FC3);
CV_Assert(rvec.size() == Size(3, 1) && rvec.type() == CV_32F);
CV_Assert(tvec.size() == Size(3, 1) && tvec.type() == CV_32F);
// Convert rotation vector into matrix
Mat rot;
Rodrigues(rvec, rot);
dst.create(src.size(), src.type());
transform_points::call(src, rot.ptr<float>(), tvec.ptr<float>(), dst, stream);
}
}
void cv::gpu::transformPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec, GpuMat& dst, Stream& stream)
{
::transformPointsCaller(src, rvec, tvec, dst, StreamAccessor::getStream(stream));
}
namespace cv { namespace gpu { namespace project_points
{
void call(const DevMem2D_<float3> src, const float* rot, const float* transl, const float* proj, DevMem2D_<float2> dst, cudaStream_t stream);
}}}
namespace
{
void projectPointsCaller(const GpuMat& src, const Mat& rvec, const Mat& tvec, const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst, cudaStream_t stream)
{
CV_Assert(src.rows == 1 && src.cols > 0 && src.type() == CV_32FC3);
CV_Assert(rvec.size() == Size(3, 1) && rvec.type() == CV_32F);
CV_Assert(tvec.size() == Size(3, 1) && tvec.type() == CV_32F);
CV_Assert(camera_mat.size() == Size(3, 3) && camera_mat.type() == CV_32F);
CV_Assert(dist_coef.empty()); // Undistortion isn't supported
// Convert rotation vector into matrix
Mat rot;
Rodrigues(rvec, rot);
dst.create(src.size(), CV_32FC2);
project_points::call(src, rot.ptr<float>(), tvec.ptr<float>(), camera_mat.ptr<float>(), dst,stream);
}
}
void cv::gpu::projectPoints(const GpuMat& src, const Mat& rvec, const Mat& tvec, const Mat& camera_mat, const Mat& dist_coef, GpuMat& dst, Stream& stream)
{
::projectPointsCaller(src, rvec, tvec, camera_mat, dist_coef, dst, StreamAccessor::getStream(stream));
}
namespace cv { namespace gpu { namespace solve_pnp_ransac
{
int maxNumIters();
void computeHypothesisScores(
const int num_hypotheses, const int num_points, const float* rot_matrices,
const float3* transl_vectors, const float3* object, const float2* image,
const float dist_threshold, int* hypothesis_scores);
}}}
namespace
{
// Selects subset_size random different points from [0, num_points - 1] range
void selectRandom(int subset_size, int num_points, vector<int>& subset)
{
subset.resize(subset_size);
for (int i = 0; i < subset_size; ++i)
{
bool was;
do
{
subset[i] = rand() % num_points;
was = false;
for (int j = 0; j < i; ++j)
if (subset[j] == subset[i])
{
was = true;
break;
}
} while (was);
}
}
// Computes rotation, translation pair for small subsets if the input data
class TransformHypothesesGenerator
{
public:
TransformHypothesesGenerator(const Mat& object_, const Mat& image_, const Mat& dist_coef_,
const Mat& camera_mat_, int num_points_, int subset_size_,
Mat rot_matrices_, Mat transl_vectors_)
: object(&object_), image(&image_), dist_coef(&dist_coef_), camera_mat(&camera_mat_),
num_points(num_points_), subset_size(subset_size_), rot_matrices(rot_matrices_),
transl_vectors(transl_vectors_) {}
void operator()(const BlockedRange& range) const
{
// Input data for generation of the current hypothesis
vector<int> subset_indices(subset_size);
Mat_<Point3f> object_subset(1, subset_size);
Mat_<Point2f> image_subset(1, subset_size);
// Current hypothesis data
Mat rot_vec(1, 3, CV_64F);
Mat rot_mat(3, 3, CV_64F);
Mat transl_vec(1, 3, CV_64F);
for (int iter = range.begin(); iter < range.end(); ++iter)
{
selectRandom(subset_size, num_points, subset_indices);
for (int i = 0; i < subset_size; ++i)
{
object_subset(0, i) = object->at<Point3f>(subset_indices[i]);
image_subset(0, i) = image->at<Point2f>(subset_indices[i]);
}
solvePnP(object_subset, image_subset, *camera_mat, *dist_coef, rot_vec, transl_vec);
// Remember translation vector
Mat transl_vec_ = transl_vectors.colRange(iter * 3, (iter + 1) * 3);
transl_vec = transl_vec.reshape(0, 1);
transl_vec.convertTo(transl_vec_, CV_32F);
// Remember rotation matrix
Rodrigues(rot_vec, rot_mat);
Mat rot_mat_ = rot_matrices.colRange(iter * 9, (iter + 1) * 9).reshape(0, 3);
rot_mat.convertTo(rot_mat_, CV_32F);
}
}
const Mat* object;
const Mat* image;
const Mat* dist_coef;
const Mat* camera_mat;
int num_points;
int subset_size;
// Hypotheses storage (global)
Mat rot_matrices;
Mat transl_vectors;
};
}
void cv::gpu::solvePnPRansac(const Mat& object, const Mat& image, const Mat& camera_mat,
const Mat& dist_coef, Mat& rvec, Mat& tvec, bool use_extrinsic_guess,
int num_iters, float max_dist, int min_inlier_count,
vector<int>* inliers)
{
CV_Assert(object.rows == 1 && object.cols > 0 && object.type() == CV_32FC3);
CV_Assert(image.rows == 1 && image.cols > 0 && image.type() == CV_32FC2);
CV_Assert(object.cols == image.cols);
CV_Assert(camera_mat.size() == Size(3, 3) && camera_mat.type() == CV_32F);
CV_Assert(!use_extrinsic_guess); // We don't support initial guess for now
CV_Assert(num_iters <= solve_pnp_ransac::maxNumIters());
const int subset_size = 4;
const int num_points = object.cols;
CV_Assert(num_points >= subset_size);
// Unapply distortion and intrinsic camera transformations
Mat eye_camera_mat = Mat::eye(3, 3, CV_32F);
Mat empty_dist_coef;
Mat image_normalized;
undistortPoints(image, image_normalized, camera_mat, dist_coef, Mat(), eye_camera_mat);
// Hypotheses storage (global)
Mat rot_matrices(1, num_iters * 9, CV_32F);
Mat transl_vectors(1, num_iters * 3, CV_32F);
// Generate set of hypotheses using small subsets of the input data
TransformHypothesesGenerator body(object, image_normalized, empty_dist_coef, eye_camera_mat,
num_points, subset_size, rot_matrices, transl_vectors);
parallel_for(BlockedRange(0, num_iters), body);
// Compute scores (i.e. number of inliers) for each hypothesis
GpuMat d_object(object);
GpuMat d_image_normalized(image_normalized);
GpuMat d_hypothesis_scores(1, num_iters, CV_32S);
solve_pnp_ransac::computeHypothesisScores(
num_iters, num_points, rot_matrices.ptr<float>(), transl_vectors.ptr<float3>(),
d_object.ptr<float3>(), d_image_normalized.ptr<float2>(), max_dist * max_dist,
d_hypothesis_scores.ptr<int>());
// Find the best hypothesis index
Point best_idx;
double best_score;
minMaxLoc(d_hypothesis_scores, NULL, &best_score, NULL, &best_idx);
int num_inliers = static_cast<int>(best_score);
// Extract the best hypothesis data
Mat rot_mat = rot_matrices.colRange(best_idx.x * 9, (best_idx.x + 1) * 9).reshape(0, 3);
Rodrigues(rot_mat, rvec);
rvec = rvec.reshape(0, 1);
tvec = transl_vectors.colRange(best_idx.x * 3, (best_idx.x + 1) * 3).clone();
tvec = tvec.reshape(0, 1);
// Build vector of inlier indices
if (inliers != NULL)
{
inliers->clear();
inliers->reserve(num_inliers);
Point3f p, p_transf;
Point2f p_proj;
const float* rot = rot_mat.ptr<float>();
const float* transl = tvec.ptr<float>();
for (int i = 0; i < num_points; ++i)
{
p = object.at<Point3f>(0, i);
p_transf.x = rot[0] * p.x + rot[1] * p.y + rot[2] * p.z + transl[0];
p_transf.y = rot[3] * p.x + rot[4] * p.y + rot[5] * p.z + transl[1];
p_transf.z = rot[6] * p.x + rot[7] * p.y + rot[8] * p.z + transl[2];
p_proj.x = p_transf.x / p_transf.z;
p_proj.y = p_transf.y / p_transf.z;
if (norm(p_proj - image_normalized.at<Point2f>(0, i)) < max_dist)
inliers->push_back(i);
}
}
}
#endif
vector<int>* inliers)
{
CV_Assert(object.rows == 1 && object.cols > 0 && object.type() == CV_32FC3);
CV_Assert(image.rows == 1 && image.cols > 0 && image.type() == CV_32FC2);
CV_Assert(object.cols == image.cols);
CV_Assert(camera_mat.size() == Size(3, 3) && camera_mat.type() == CV_32F);
CV_Assert(!use_extrinsic_guess); // We don't support initial guess for now
CV_Assert(num_iters <= solve_pnp_ransac::maxNumIters());
const int subset_size = 4;
const int num_points = object.cols;
CV_Assert(num_points >= subset_size);
// Unapply distortion and intrinsic camera transformations
Mat eye_camera_mat = Mat::eye(3, 3, CV_32F);
Mat empty_dist_coef;
Mat image_normalized;
undistortPoints(image, image_normalized, camera_mat, dist_coef, Mat(), eye_camera_mat);
// Hypotheses storage (global)
Mat rot_matrices(1, num_iters * 9, CV_32F);
Mat transl_vectors(1, num_iters * 3, CV_32F);
// Generate set of hypotheses using small subsets of the input data
TransformHypothesesGenerator body(object, image_normalized, empty_dist_coef, eye_camera_mat,
num_points, subset_size, rot_matrices, transl_vectors);
parallel_for(BlockedRange(0, num_iters), body);
// Compute scores (i.e. number of inliers) for each hypothesis
GpuMat d_object(object);
GpuMat d_image_normalized(image_normalized);
GpuMat d_hypothesis_scores(1, num_iters, CV_32S);
solve_pnp_ransac::computeHypothesisScores(
num_iters, num_points, rot_matrices.ptr<float>(), transl_vectors.ptr<float3>(),
d_object.ptr<float3>(), d_image_normalized.ptr<float2>(), max_dist * max_dist,
d_hypothesis_scores.ptr<int>());
// Find the best hypothesis index
Point best_idx;
double best_score;
minMaxLoc(d_hypothesis_scores, NULL, &best_score, NULL, &best_idx);
int num_inliers = static_cast<int>(best_score);
// Extract the best hypothesis data
Mat rot_mat = rot_matrices.colRange(best_idx.x * 9, (best_idx.x + 1) * 9).reshape(0, 3);
Rodrigues(rot_mat, rvec);
rvec = rvec.reshape(0, 1);
tvec = transl_vectors.colRange(best_idx.x * 3, (best_idx.x + 1) * 3).clone();
tvec = tvec.reshape(0, 1);
// Build vector of inlier indices
if (inliers != NULL)
{
inliers->clear();
inliers->reserve(num_inliers);
Point3f p, p_transf;
Point2f p_proj;
const float* rot = rot_mat.ptr<float>();
const float* transl = tvec.ptr<float>();
for (int i = 0; i < num_points; ++i)
{
p = object.at<Point3f>(0, i);
p_transf.x = rot[0] * p.x + rot[1] * p.y + rot[2] * p.z + transl[0];
p_transf.y = rot[3] * p.x + rot[4] * p.y + rot[5] * p.z + transl[1];
p_transf.z = rot[6] * p.x + rot[7] * p.y + rot[8] * p.z + transl[2];
p_proj.x = p_transf.x / p_transf.z;
p_proj.y = p_transf.y / p_transf.z;
if (norm(p_proj - image_normalized.at<Point2f>(0, i)) < max_dist)
inliers->push_back(i);
}
}
}
#endif

4
modules/gpu/src/cuda/brute_force_matcher.cu
View File

@ -646,9 +646,9 @@ namespace cv { namespace gpu { namespace bfmatcher
matchCached_caller<16, 16, 64, true, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance, stream);
else if (queryDescs.cols < 128)
matchCached_caller<16, 16, 128, false, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance, stream);
else if (queryDescs.cols == 128)
else if (queryDescs.cols == 128 && cc_12)
matchCached_caller<16, 16, 128, true, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance, stream);
else if (queryDescs.cols < 256)
else if (queryDescs.cols < 256 && cc_12)
matchCached_caller<16, 16, 256, false, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance, stream);
else if (queryDescs.cols == 256 && cc_12)
matchCached_caller<16, 16, 256, true, Dist>(queryDescs, train, mask, trainIdx, imgIdx, distance, stream);

4
modules/gpu/src/cuda/mathfunc.cu
View File

@ -82,7 +82,9 @@ namespace cv { namespace gpu { namespace mathfunc
{
static __device__ __forceinline__ void calc(int x, int y, float x_data, float y_data, float* dst, size_t dst_step, float scale)
{
dst[y * dst_step + x] = scale * atan2f(y_data, x_data);
float angle = atan2f(y_data, x_data);
angle += (angle < 0) * 2.0 * CV_PI;
dst[y * dst_step + x] = scale * angle;
}
};
template <typename Mag, typename Angle>

36
modules/gpu/src/element_operations.cpp
View File

@ -211,22 +211,42 @@ void cv::gpu::subtract(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream&
void cv::gpu::multiply(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream& stream)
{
typedef void (*caller_t)(const GpuMat& src, const Scalar& sc, GpuMat& dst, cudaStream_t stream);
static const caller_t callers[] = {0, NppArithmScalar<1, nppiMulC_32f_C1R>::calc, NppArithmScalar<2, nppiMulC_32fc_C1R>::calc};
CV_Assert(src.type() == CV_32FC1);
CV_Assert(src.type() == CV_32FC1 || src.type() == CV_32FC2);
dst.create(src.size(), src.type());
callers[src.channels()](src, sc, dst, StreamAccessor::getStream(stream));
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
cudaStream_t cudaStream = StreamAccessor::getStream(stream);
NppStreamHandler h(cudaStream);
nppSafeCall( nppiMulC_32f_C1R(src.ptr<Npp32f>(), src.step, (Npp32f)sc[0], dst.ptr<Npp32f>(), dst.step, sz) );
if (cudaStream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}
void cv::gpu::divide(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream& stream)
{
typedef void (*caller_t)(const GpuMat& src, const Scalar& sc, GpuMat& dst, cudaStream_t stream);
static const caller_t callers[] = {0, NppArithmScalar<1, nppiDivC_32f_C1R>::calc, NppArithmScalar<2, nppiDivC_32fc_C1R>::calc};
CV_Assert(src.type() == CV_32FC1);
CV_Assert(src.type() == CV_32FC1 || src.type() == CV_32FC2);
dst.create(src.size(), src.type());
callers[src.channels()](src, sc, dst, StreamAccessor::getStream(stream));
NppiSize sz;
sz.width = src.cols;
sz.height = src.rows;
cudaStream_t cudaStream = StreamAccessor::getStream(stream);
NppStreamHandler h(cudaStream);
nppSafeCall( nppiDivC_32f_C1R(src.ptr<Npp32f>(), src.step, (Npp32f)sc[0], dst.ptr<Npp32f>(), dst.step, sz) );
if (cudaStream == 0)
cudaSafeCall( cudaDeviceSynchronize() );
}

44
modules/gpu/test/nvidia/NCVAutoTestLister.hpp
View File

@ -15,15 +15,22 @@
#include "NCVTest.hpp"
enum OutputLevel
{
OutputLevelNone,
OutputLevelCompact,
OutputLevelFull
};
class NCVAutoTestLister
{
public:
NCVAutoTestLister(std::string testSuiteName, NcvBool bStopOnFirstFail=false, NcvBool bCompactOutput=true)
NCVAutoTestLister(std::string testSuiteName, OutputLevel outputLevel = OutputLevelCompact, NcvBool bStopOnFirstFail=false)
:
testSuiteName(testSuiteName),
bStopOnFirstFail(bStopOnFirstFail),
bCompactOutput(bCompactOutput)
outputLevel(outputLevel),
bStopOnFirstFail(bStopOnFirstFail)
{
}
@ -38,7 +45,7 @@ public:
Ncv32u nFailed = 0;
Ncv32u nFailedMem = 0;
if (bCompactOutput)
if (outputLevel == OutputLevelCompact)
{
printf("Test suite '%s' with %d tests\n",
testSuiteName.c_str(),
@ -52,7 +59,7 @@ public:
NCVTestReport curReport;
bool res = curTest.executeTest(curReport);
if (!bCompactOutput)
if (outputLevel == OutputLevelFull)
{
printf("Test %3i %16s; Consumed mem GPU = %8d, CPU = %8d; %s\n",
i,
@ -65,7 +72,7 @@ public:
if (res)
{
nPassed++;
if (bCompactOutput)
if (outputLevel == OutputLevelCompact)
{
printf(".");
}
@ -75,7 +82,7 @@ public:
if (!curReport.statsText["rcode"].compare("FAILED"))
{
nFailed++;
if (bCompactOutput)
if (outputLevel == OutputLevelCompact)
{
printf("x");
}
@ -87,7 +94,7 @@ public:
else
{
nFailedMem++;
if (bCompactOutput)
if (outputLevel == OutputLevelCompact)
{
printf("m");
}
@ -95,17 +102,20 @@ public:
}
fflush(stdout);
}
if (bCompactOutput)
if (outputLevel == OutputLevelCompact)
{
printf("\n");
}
printf("Test suite '%s' complete: %d total, %d passed, %d memory errors, %d failed\n\n",
testSuiteName.c_str(),
(int)(this->tests.size()),
nPassed,
nFailedMem,
nFailed);
if (outputLevel != OutputLevelNone)
{
printf("Test suite '%s' complete: %d total, %d passed, %d memory errors, %d failed\n\n",
testSuiteName.c_str(),
(int)(this->tests.size()),
nPassed,
nFailedMem,
nFailed);
}
bool passed = nFailed == 0 && nFailedMem == 0;
return passed;
@ -121,9 +131,9 @@ public:
private:
NcvBool bStopOnFirstFail;
NcvBool bCompactOutput;
std::string testSuiteName;
OutputLevel outputLevel;
NcvBool bStopOnFirstFail;
std::vector<INCVTest *> tests;
};

182
modules/gpu/test/nvidia/main_nvidia.cpp
View File

@ -288,70 +288,162 @@ static void devNullOutput(const char *msg)
}
bool main_nvidia(const std::string& test_data_path)
bool nvidia_NPPST_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
printf("Testing NVIDIA Computer Vision SDK\n");
printf("==================================\n");
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerII("NPPST Integral Image" );//,,true, false);
NCVAutoTestLister testListerSII("NPPST Squared Integral Image" );//,,true, false);
NCVAutoTestLister testListerRStdDev("NPPST RectStdDev" );//,,true, false);
NCVAutoTestLister testListerResize("NPPST Resize" );//,,true, false);
NCVAutoTestLister testListerNPPSTVectorOperations("NPPST Vector Operations" );//,,true, false);
NCVAutoTestLister testListerTranspose("NPPST Transpose" );//,,true, false);
NCVAutoTestLister testListerVectorOperations("Vector Operations" );//,,true, false);
NCVAutoTestLister testListerHaarLoader("Haar Cascade Loader" );//,,true, false);
NCVAutoTestLister testListerHaarAppl("Haar Cascade Application" );//,,true, false);
NCVAutoTestLister testListerHypFiltration("Hypotheses Filtration" );//,,true, false);
NCVAutoTestLister testListerVisualize("Visualization" );//,,true, false);
printf("Initializing data source providers\n");
NCVTestSourceProvider<Ncv32u> testSrcRandom_32u(2010, 0, 0xFFFFFFFF, 4096, 4096);
NCVAutoTestLister testListerII("NPPST Integral Image", outputLevel);
NCVTestSourceProvider<Ncv8u> testSrcRandom_8u(2010, 0, 255, 4096, 4096);
NCVTestSourceProvider<Ncv64u> testSrcRandom_64u(2010, 0, -1, 4096, 4096);
NCVTestSourceProvider<Ncv8u> testSrcFacesVGA_8u(path + "group_1_640x480_VGA.pgm");
NCVTestSourceProvider<Ncv32f> testSrcRandom_32f(2010, -1.0f, 1.0f, 4096, 4096);
printf("Generating NPPST test suites\n");
generateIntegralTests<Ncv8u, Ncv32u>(testListerII, testSrcRandom_8u, 4096, 4096);
generateIntegralTests<Ncv32f, Ncv32f>(testListerII, testSrcRandom_32f, 4096, 4096);
return testListerII.invoke();
}
bool nvidia_NPPST_Squared_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerSII("NPPST Squared Integral Image", outputLevel);
NCVTestSourceProvider<Ncv8u> testSrcRandom_8u(2010, 0, 255, 4096, 4096);
generateSquaredIntegralTests(testListerSII, testSrcRandom_8u, 4096, 4096);
return testListerSII.invoke();
}
bool nvidia_NPPST_RectStdDev(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerRStdDev("NPPST RectStdDev", outputLevel);
NCVTestSourceProvider<Ncv8u> testSrcRandom_8u(2010, 0, 255, 4096, 4096);
generateRectStdDevTests(testListerRStdDev, testSrcRandom_8u, 4096, 4096);
return testListerRStdDev.invoke();
}
bool nvidia_NPPST_Resize(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerResize("NPPST Resize", outputLevel);
NCVTestSourceProvider<Ncv32u> testSrcRandom_32u(2010, 0, 0xFFFFFFFF, 4096, 4096);
NCVTestSourceProvider<Ncv64u> testSrcRandom_64u(2010, 0, -1, 4096, 4096);
generateResizeTests(testListerResize, testSrcRandom_32u);
generateResizeTests(testListerResize, testSrcRandom_64u);
return testListerResize.invoke();
}
bool nvidia_NPPST_Vector_Operations(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerNPPSTVectorOperations("NPPST Vector Operations", outputLevel);
NCVTestSourceProvider<Ncv32u> testSrcRandom_32u(2010, 0, 0xFFFFFFFF, 4096, 4096);
generateNPPSTVectorTests(testListerNPPSTVectorOperations, testSrcRandom_32u, 4096*4096);
return testListerNPPSTVectorOperations.invoke();
}
bool nvidia_NPPST_Transpose(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerTranspose("NPPST Transpose", outputLevel);
NCVTestSourceProvider<Ncv32u> testSrcRandom_32u(2010, 0, 0xFFFFFFFF, 4096, 4096);
NCVTestSourceProvider<Ncv64u> testSrcRandom_64u(2010, 0, -1, 4096, 4096);
generateTransposeTests(testListerTranspose, testSrcRandom_32u);
generateTransposeTests(testListerTranspose, testSrcRandom_64u);
printf("Generating NCV test suites\n");
generateDrawRectsTests(testListerVisualize, testSrcRandom_8u, testSrcRandom_32u, 4096, 4096);
generateDrawRectsTests(testListerVisualize, testSrcRandom_32u, testSrcRandom_32u, 4096, 4096);
return testListerTranspose.invoke();
}
bool nvidia_NCV_Vector_Operations(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerVectorOperations("Vector Operations", outputLevel);
NCVTestSourceProvider<Ncv32u> testSrcRandom_32u(2010, 0, 0xFFFFFFFF, 4096, 4096);
generateVectorTests(testListerVectorOperations, testSrcRandom_32u, 4096*4096);
generateHypothesesFiltrationTests(testListerHypFiltration, testSrcRandom_32u, 1024);
return testListerVectorOperations.invoke();
}
bool nvidia_NCV_Haar_Cascade_Loader(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerHaarLoader("Haar Cascade Loader", outputLevel);
generateHaarLoaderTests(testListerHaarLoader);
return testListerHaarLoader.invoke();
}
bool nvidia_NCV_Haar_Cascade_Application(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerHaarAppl("Haar Cascade Application", outputLevel);
NCVTestSourceProvider<Ncv8u> testSrcFacesVGA_8u(path + "group_1_640x480_VGA.pgm");
generateHaarApplicationTests(testListerHaarAppl, testSrcFacesVGA_8u, 1280, 720);
// Indicate if at least one test failed
bool passed = true;
// Invoke all tests
passed &= testListerII.invoke();
passed &= testListerSII.invoke();
passed &= testListerRStdDev.invoke();
passed &= testListerResize.invoke();
passed &= testListerNPPSTVectorOperations.invoke();
passed &= testListerTranspose.invoke();
passed &= testListerVisualize.invoke();
passed &= testListerVectorOperations.invoke();
passed &= testListerHypFiltration.invoke();
passed &= testListerHaarLoader.invoke();
passed &= testListerHaarAppl.invoke();
return passed;
return testListerHaarAppl.invoke();
}
bool nvidia_NCV_Hypotheses_Filtration(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerHypFiltration("Hypotheses Filtration", outputLevel);
NCVTestSourceProvider<Ncv32u> testSrcRandom_32u(2010, 0, 0xFFFFFFFF, 4096, 4096);
generateHypothesesFiltrationTests(testListerHypFiltration, testSrcRandom_32u, 1024);
return testListerHypFiltration.invoke();
}
bool nvidia_NCV_Visualization(const std::string& test_data_path, OutputLevel outputLevel)
{
path = test_data_path;
ncvSetDebugOutputHandler(devNullOutput);
NCVAutoTestLister testListerVisualize("Visualization", outputLevel);
NCVTestSourceProvider<Ncv8u> testSrcRandom_8u(2010, 0, 255, 4096, 4096);
NCVTestSourceProvider<Ncv32u> testSrcRandom_32u(2010, 0, 0xFFFFFFFF, 4096, 4096);
generateDrawRectsTests(testListerVisualize, testSrcRandom_8u, testSrcRandom_32u, 4096, 4096);
generateDrawRectsTests(testListerVisualize, testSrcRandom_32u, testSrcRandom_32u, 4096, 4096);
return testListerVisualize.invoke();
}

2153
modules/gpu/test/test_arithm.cpp
View File

File diff suppressed because it is too large Load Diff

235
modules/gpu/test/test_bitwise_oper.cpp
View File

@ -1,235 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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 <iostream>
#include <limits>
#include "test_precomp.hpp"
#define CHECK(pred, err) if (!(pred)) { \
ts->printf(cvtest::TS::CONSOLE, "Fail: \"%s\" at line: %d\n", #pred, __LINE__); \
ts->set_failed_test_info(err); \
return; }
using namespace cv;
using namespace std;
struct CV_GpuBitwiseTest: public cvtest::BaseTest
{
CV_GpuBitwiseTest() {}
void run(int)
{
int rows, cols;
bool double_ok = gpu::TargetArchs::builtWith(gpu::NATIVE_DOUBLE) &&
gpu::DeviceInfo().supports(gpu::NATIVE_DOUBLE);
int depth_end = double_ok ? CV_64F : CV_32F;
for (int depth = CV_8U; depth <= depth_end; ++depth)
for (int cn = 1; cn <= 4; ++cn)
for (int attempt = 0; attempt < 3; ++attempt)
{
rows = 1 + rand() % 100;
cols = 1 + rand() % 100;
test_bitwise_not(rows, cols, CV_MAKETYPE(depth, cn));
test_bitwise_or(rows, cols, CV_MAKETYPE(depth, cn));
test_bitwise_and(rows, cols, CV_MAKETYPE(depth, cn));
test_bitwise_xor(rows, cols, CV_MAKETYPE(depth, cn));
}
}
void test_bitwise_not(int rows, int cols, int type)
{
Mat src(rows, cols, type);
RNG rng;
for (int i = 0; i < src.rows; ++i)
{
Mat row(1, src.cols * src.elemSize(), CV_8U, src.ptr(i));
rng.fill(row, RNG::UNIFORM, Scalar(0), Scalar(255));
}
Mat dst_gold = ~src;
gpu::GpuMat mask(src.size(), CV_8U);
mask.setTo(Scalar(1));
gpu::GpuMat dst;
gpu::bitwise_not(gpu::GpuMat(src), dst);
CHECK(dst_gold.size() == dst.size(), cvtest::TS::FAIL_INVALID_OUTPUT);
CHECK(dst_gold.type() == dst.type(), cvtest::TS::FAIL_INVALID_OUTPUT);
Mat dsth(dst);
for (int i = 0; i < dst_gold.rows; ++i)
CHECK(memcmp(dst_gold.ptr(i), dsth.ptr(i), dst_gold.cols * dst_gold.elemSize()) == 0, cvtest::TS::FAIL_INVALID_OUTPUT);
dst.setTo(Scalar::all(0));
gpu::bitwise_not(gpu::GpuMat(src), dst, mask);
CHECK(dst_gold.size() == dst.size(), cvtest::TS::FAIL_INVALID_OUTPUT);
CHECK(dst_gold.type() == dst.type(), cvtest::TS::FAIL_INVALID_OUTPUT);
dsth = dst;
for (int i = 0; i < dst_gold.rows; ++i)
CHECK(memcmp(dst_gold.ptr(i), dsth.ptr(i), dst_gold.cols * dst_gold.elemSize()) == 0, cvtest::TS::FAIL_INVALID_OUTPUT)
}
void test_bitwise_or(int rows, int cols, int type)
{
Mat src1(rows, cols, type);
Mat src2(rows, cols, type);
RNG rng;
for (int i = 0; i < src1.rows; ++i)
{
Mat row1(1, src1.cols * src1.elemSize(), CV_8U, src1.ptr(i));
rng.fill(row1, RNG::UNIFORM, Scalar(0), Scalar(255));
Mat row2(1, src2.cols * src2.elemSize(), CV_8U, src2.ptr(i));
rng.fill(row2, RNG::UNIFORM, Scalar(0), Scalar(255));
}
Mat dst_gold = src1 | src2;
gpu::GpuMat dst = gpu::GpuMat(src1) | gpu::GpuMat(src2);
CHECK(dst_gold.size() == dst.size(), cvtest::TS::FAIL_INVALID_OUTPUT);
CHECK(dst_gold.type() == dst.type(), cvtest::TS::FAIL_INVALID_OUTPUT);
Mat dsth(dst);
for (int i = 0; i < dst_gold.rows; ++i)
CHECK(memcmp(dst_gold.ptr(i), dsth.ptr(i), dst_gold.cols * dst_gold.elemSize()) == 0, cvtest::TS::FAIL_INVALID_OUTPUT)
Mat mask(src1.size(), CV_8U);
randu(mask, Scalar(0), Scalar(255));
Mat dst_gold2(dst_gold.size(), dst_gold.type()); dst_gold2.setTo(Scalar::all(0));
gpu::GpuMat dst2(dst.size(), dst.type()); dst2.setTo(Scalar::all(0));
bitwise_or(src1, src2, dst_gold2, mask);
gpu::bitwise_or(gpu::GpuMat(src1), gpu::GpuMat(src2), dst2, gpu::GpuMat(mask));
CHECK(dst_gold2.size() == dst2.size(), cvtest::TS::FAIL_INVALID_OUTPUT);
CHECK(dst_gold2.type() == dst2.type(), cvtest::TS::FAIL_INVALID_OUTPUT);
dsth = dst2;
for (int i = 0; i < dst_gold.rows; ++i)
CHECK(memcmp(dst_gold2.ptr(i), dsth.ptr(i), dst_gold2.cols * dst_gold2.elemSize()) == 0, cvtest::TS::FAIL_INVALID_OUTPUT)
}
void test_bitwise_and(int rows, int cols, int type)
{
Mat src1(rows, cols, type);
Mat src2(rows, cols, type);
RNG rng;
for (int i = 0; i < src1.rows; ++i)
{
Mat row1(1, src1.cols * src1.elemSize(), CV_8U, src1.ptr(i));
rng.fill(row1, RNG::UNIFORM, Scalar(0), Scalar(255));
Mat row2(1, src2.cols * src2.elemSize(), CV_8U, src2.ptr(i));
rng.fill(row2, RNG::UNIFORM, Scalar(0), Scalar(255));
}
Mat dst_gold = src1 & src2;
gpu::GpuMat dst = gpu::GpuMat(src1) & gpu::GpuMat(src2);
CHECK(dst_gold.size() == dst.size(), cvtest::TS::FAIL_INVALID_OUTPUT);
CHECK(dst_gold.type() == dst.type(), cvtest::TS::FAIL_INVALID_OUTPUT);
Mat dsth(dst);
for (int i = 0; i < dst_gold.rows; ++i)
CHECK(memcmp(dst_gold.ptr(i), dsth.ptr(i), dst_gold.cols * dst_gold.elemSize()) == 0, cvtest::TS::FAIL_INVALID_OUTPUT)
Mat mask(src1.size(), CV_8U);
randu(mask, Scalar(0), Scalar(255));
Mat dst_gold2(dst_gold.size(), dst_gold.type()); dst_gold2.setTo(Scalar::all(0));
gpu::GpuMat dst2(dst.size(), dst.type()); dst2.setTo(Scalar::all(0));
bitwise_and(src1, src2, dst_gold2, mask);
gpu::bitwise_and(gpu::GpuMat(src1), gpu::GpuMat(src2), dst2, gpu::GpuMat(mask));
CHECK(dst_gold2.size() == dst2.size(), cvtest::TS::FAIL_INVALID_OUTPUT);
CHECK(dst_gold2.type() == dst2.type(), cvtest::TS::FAIL_INVALID_OUTPUT);
dsth = dst2;
for (int i = 0; i < dst_gold.rows; ++i)
CHECK(memcmp(dst_gold2.ptr(i), dsth.ptr(i), dst_gold2.cols * dst_gold2.elemSize()) == 0, cvtest::TS::FAIL_INVALID_OUTPUT)
}
void test_bitwise_xor(int rows, int cols, int type)
{
Mat src1(rows, cols, type);
Mat src2(rows, cols, type);
RNG rng;
for (int i = 0; i < src1.rows; ++i)
{
Mat row1(1, src1.cols * src1.elemSize(), CV_8U, src1.ptr(i));
rng.fill(row1, RNG::UNIFORM, Scalar(0), Scalar(255));
Mat row2(1, src2.cols * src2.elemSize(), CV_8U, src2.ptr(i));
rng.fill(row2, RNG::UNIFORM, Scalar(0), Scalar(255));
}
Mat dst_gold = src1 ^ src2;
gpu::GpuMat dst = gpu::GpuMat(src1) ^ gpu::GpuMat(src2);
CHECK(dst_gold.size() == dst.size(), cvtest::TS::FAIL_INVALID_OUTPUT);
CHECK(dst_gold.type() == dst.type(), cvtest::TS::FAIL_INVALID_OUTPUT);
Mat dsth(dst);
for (int i = 0; i < dst_gold.rows; ++i)
CHECK(memcmp(dst_gold.ptr(i), dsth.ptr(i), dst_gold.cols * dst_gold.elemSize()) == 0, cvtest::TS::FAIL_INVALID_OUTPUT)
Mat mask(src1.size(), CV_8U);
randu(mask, Scalar(0), Scalar(255));
Mat dst_gold2(dst_gold.size(), dst_gold.type()); dst_gold2.setTo(Scalar::all(0));
gpu::GpuMat dst2(dst.size(), dst.type()); dst2.setTo(Scalar::all(0));
bitwise_xor(src1, src2, dst_gold2, mask);
gpu::bitwise_xor(gpu::GpuMat(src1), gpu::GpuMat(src2), dst2, gpu::GpuMat(mask));
CHECK(dst_gold2.size() == dst2.size(), cvtest::TS::FAIL_INVALID_OUTPUT);
CHECK(dst_gold2.type() == dst2.type(), cvtest::TS::FAIL_INVALID_OUTPUT);
dsth = dst2;
for (int i = 0; i < dst_gold.rows; ++i)
CHECK(memcmp(dst_gold2.ptr(i), dsth.ptr(i), dst_gold2.cols * dst_gold2.elemSize()) == 0, cvtest::TS::FAIL_INVALID_OUTPUT)
}
};
TEST(BitwiseOperations, accuracy) { CV_GpuBitwiseTest test; test.safe_run(); }

100
modules/gpu/test/test_blend.cpp
View File

@ -1,100 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
using namespace std;
using namespace cv;
using namespace cv::gpu;
TEST(blendLinear, accuracy_on_8U)
{
RNG& rng = cvtest::TS::ptr()->get_rng();
Size size(200 + cvtest::randInt(rng) % 1000,
200 + cvtest::randInt(rng) % 1000);
for (int cn = 1; cn <= 4; ++cn)
{
Mat img1 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_8U, cn), 0, 255, false);
Mat img2 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_8U, cn), 0, 255, false);
Mat weights1 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);
Mat weights2 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);
Mat result_gold(size, CV_MAKE_TYPE(CV_8U, cn));
for (int y = 0; y < size.height; ++y)
for (int x = 0; x < size.width * cn; ++x)
{
float w1 = weights1.at<float>(y, x / cn);
float w2 = weights2.at<float>(y, x / cn);
result_gold.at<uchar>(y, x) = static_cast<uchar>(
(img1.at<uchar>(y, x) * w1 + img2.at<uchar>(y, x) * w2) / (w1 + w2 + 1e-5f));
}
GpuMat d_result;
blendLinear(GpuMat(img1), GpuMat(img2), GpuMat(weights1), GpuMat(weights2), d_result);
ASSERT_LE(cvtest::norm(result_gold, Mat(d_result), NORM_INF), 1)
<< "rows=" << size.height << ", cols=" << size.width << ", cn=" << cn;
}
}
TEST(blendLinear, accuracy_on_32F)
{
RNG& rng = cvtest::TS::ptr()->get_rng();
Size size(200 + cvtest::randInt(rng) % 1000,
200 + cvtest::randInt(rng) % 1000);
for (int cn = 1; cn <= 4; ++cn)
{
Mat img1 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_32F, cn), 0, 1, false);
Mat img2 = cvtest::randomMat(rng, size, CV_MAKE_TYPE(CV_32F, cn), 0, 1, false);
Mat weights1 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);
Mat weights2 = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);
Mat result_gold(size, CV_MAKE_TYPE(CV_32F, cn));
for (int y = 0; y < size.height; ++y)
for (int x = 0; x < size.width * cn; ++x)
{
float w1 = weights1.at<float>(y, x / cn);
float w2 = weights2.at<float>(y, x / cn);
result_gold.at<float>(y, x) =
(img1.at<float>(y, x) * w1 + img2.at<float>(y, x) * w2) / (w1 + w2 + 1e-5f);
}
GpuMat d_result;
blendLinear(GpuMat(img1), GpuMat(img2), GpuMat(weights1), GpuMat(weights2), d_result);
ASSERT_LE(cvtest::norm(result_gold, Mat(d_result), NORM_INF), 1e-3)
<< "rows=" << size.height << ", cols=" << size.width << ", cn=" << cn;
}
}

522
modules/gpu/test/test_brute_force_matcher.cpp
View File

@ -1,522 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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 <algorithm>
#include <iterator>
using namespace cv;
using namespace cv::gpu;
using namespace std;
class CV_GpuBruteForceMatcherTest : public cvtest::BaseTest
{
public:
CV_GpuBruteForceMatcherTest()
{
}
protected:
virtual void run(int);
void emptyDataTest();
void dataTest(int dim);
void generateData(GpuMat& query, GpuMat& train, int dim);
void matchTest(const GpuMat& query, const GpuMat& train);
void knnMatchTest(const GpuMat& query, const GpuMat& train);
void radiusMatchTest(const GpuMat& query, const GpuMat& train);
private:
BruteForceMatcher_GPU< L2<float> > dmatcher;
static const int queryDescCount = 300; // must be even number because we split train data in some cases in two
static const int countFactor = 4; // do not change it
};
void CV_GpuBruteForceMatcherTest::emptyDataTest()
{
GpuMat queryDescriptors, trainDescriptors, mask;
vector<GpuMat> trainDescriptorCollection, masks;
vector<DMatch> matches;
vector< vector<DMatch> > vmatches;
try
{
dmatcher.match(queryDescriptors, trainDescriptors, matches, mask);
}
catch(...)
{
ts->printf( cvtest::TS::LOG, "match() on empty descriptors must not generate exception (1).\n" );
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
try
{
dmatcher.knnMatch(queryDescriptors, trainDescriptors, vmatches, 2, mask);
}
catch(...)
{
ts->printf( cvtest::TS::LOG, "knnMatch() on empty descriptors must not generate exception (1).\n" );
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
try
{
dmatcher.radiusMatch(queryDescriptors, trainDescriptors, vmatches, 10.f, mask);
}
catch(...)
{
ts->printf( cvtest::TS::LOG, "radiusMatch() on empty descriptors must not generate exception (1).\n" );
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
try
{
dmatcher.add(trainDescriptorCollection);
}
catch(...)
{
ts->printf( cvtest::TS::LOG, "add() on empty descriptors must not generate exception.\n" );
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
try
{
dmatcher.match(queryDescriptors, matches, masks);
}
catch(...)
{
ts->printf( cvtest::TS::LOG, "match() on empty descriptors must not generate exception (2).\n" );
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
try
{
dmatcher.knnMatch(queryDescriptors, vmatches, 2, masks);
}
catch(...)
{
ts->printf( cvtest::TS::LOG, "knnMatch() on empty descriptors must not generate exception (2).\n" );
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
try
{
dmatcher.radiusMatch( queryDescriptors, vmatches, 10.f, masks );
}
catch(...)
{
ts->printf( cvtest::TS::LOG, "radiusMatch() on empty descriptors must not generate exception (2).\n" );
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
}
void CV_GpuBruteForceMatcherTest::generateData( GpuMat& queryGPU, GpuMat& trainGPU, int dim )
{
Mat query, train;
RNG& rng = ts->get_rng();
// Generate query descriptors randomly.
// Descriptor vector elements are integer values.
Mat buf( queryDescCount, dim, CV_32SC1 );
rng.fill( buf, RNG::UNIFORM, Scalar::all(0), Scalar(3) );
buf.convertTo( query, CV_32FC1 );
// Generate train decriptors as follows:
// copy each query descriptor to train set countFactor times
// and perturb some one element of the copied descriptors in
// in ascending order. General boundaries of the perturbation
// are (0.f, 1.f).
train.create( query.rows*countFactor, query.cols, CV_32FC1 );
float step = 1.f / countFactor;
for( int qIdx = 0; qIdx < query.rows; qIdx++ )
{
Mat queryDescriptor = query.row(qIdx);
for( int c = 0; c < countFactor; c++ )
{
int tIdx = qIdx * countFactor + c;
Mat trainDescriptor = train.row(tIdx);
queryDescriptor.copyTo( trainDescriptor );
int elem = rng(dim);
float diff = rng.uniform( step*c, step*(c+1) );
trainDescriptor.at<float>(0, elem) += diff;
}
}
queryGPU.upload(query);
trainGPU.upload(train);
}
void CV_GpuBruteForceMatcherTest::matchTest( const GpuMat& query, const GpuMat& train )
{
dmatcher.clear();
// test const version of match()
{
vector<DMatch> matches;
dmatcher.match( query, train, matches );
if( (int)matches.size() != queryDescCount )
{
ts->printf(cvtest::TS::LOG, "Incorrect matches count while test match() function (1).\n");
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
else
{
int badCount = 0;
for( size_t i = 0; i < matches.size(); i++ )
{
DMatch match = matches[i];
if( (match.queryIdx != (int)i) || (match.trainIdx != (int)i*countFactor) || (match.imgIdx != 0) )
badCount++;
}
if (badCount > 0)
{
ts->printf( cvtest::TS::LOG, "%f - too large bad matches part while test match() function (1).\n",
(float)badCount/(float)queryDescCount );
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
}
}
// test version of match() with add()
{
vector<DMatch> matches;
// make add() twice to test such case
dmatcher.add( vector<GpuMat>(1,train.rowRange(0, train.rows/2)) );
dmatcher.add( vector<GpuMat>(1,train.rowRange(train.rows/2, train.rows)) );
// prepare masks (make first nearest match illegal)
vector<GpuMat> masks(2);
for(int mi = 0; mi < 2; mi++ )
{
masks[mi] = GpuMat(query.rows, train.rows/2, CV_8UC1, Scalar::all(1));
for( int di = 0; di < queryDescCount/2; di++ )
masks[mi].col(di*countFactor).setTo(Scalar::all(0));
}
dmatcher.match( query, matches, masks );
if( (int)matches.size() != queryDescCount )
{
ts->printf(cvtest::TS::LOG, "Incorrect matches count while test match() function (2).\n");
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
else
{
int badCount = 0;
for( size_t i = 0; i < matches.size(); i++ )
{
DMatch match = matches[i];
int shift = dmatcher.isMaskSupported() ? 1 : 0;
{
if( i < queryDescCount/2 )
{
if( (match.queryIdx != (int)i) || (match.trainIdx != (int)i*countFactor + shift) || (match.imgIdx != 0) )
badCount++;
}
else
{
if( (match.queryIdx != (int)i) || (match.trainIdx != ((int)i-queryDescCount/2)*countFactor + shift) || (match.imgIdx != 1) )
badCount++;
}
}
}
if (badCount > 0)
{
ts->printf( cvtest::TS::LOG, "%f - too large bad matches part while test match() function (2).\n",
(float)badCount/(float)queryDescCount );
ts->set_failed_test_info( cvtest::TS::FAIL_BAD_ACCURACY );
}
}
}
}
void CV_GpuBruteForceMatcherTest::knnMatchTest( const GpuMat& query, const GpuMat& train )
{
dmatcher.clear();
// test const version of knnMatch()
{
const int knn = 3;
vector< vector<DMatch> > matches;
dmatcher.knnMatch( query, train, matches, knn );
if( (int)matches.size() != queryDescCount )
{
ts->printf(cvtest::TS::LOG, "Incorrect matches count while test knnMatch() function (1).\n");
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
else
{
int badCount = 0;
for( size_t i = 0; i < matches.size(); i++ )
{
if( (int)matches[i].size() != knn )
badCount++;
else
{
int localBadCount = 0;
for( int k = 0; k < knn; k++ )
{
DMatch match = matches[i][k];
if( (match.queryIdx != (int)i) || (match.trainIdx != (int)i*countFactor+k) || (match.imgIdx != 0) )
localBadCount++;
}
badCount += localBadCount > 0 ? 1 : 0;
}
}
if (badCount > 0)
{
ts->printf( cvtest::TS::LOG, "%f - too large bad matches part while test knnMatch() function (1).\n",
(float)badCount/(float)queryDescCount );
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
}
}
// test version of knnMatch() with add()
{
const int knn = 2;
vector<vector<DMatch> > matches;
// make add() twice to test such case
dmatcher.add( vector<GpuMat>(1,train.rowRange(0, train.rows/2)) );
dmatcher.add( vector<GpuMat>(1,train.rowRange(train.rows/2, train.rows)) );
// prepare masks (make first nearest match illegal)
vector<GpuMat> masks(2);
for(int mi = 0; mi < 2; mi++ )
{
masks[mi] = GpuMat(query.rows, train.rows/2, CV_8UC1, Scalar::all(1));
for( int di = 0; di < queryDescCount/2; di++ )
masks[mi].col(di*countFactor).setTo(Scalar::all(0));
}
dmatcher.knnMatch( query, matches, knn, masks );
if( (int)matches.size() != queryDescCount )
{
ts->printf(cvtest::TS::LOG, "Incorrect matches count while test knnMatch() function (2).\n");
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
else
{
int badCount = 0;
int shift = dmatcher.isMaskSupported() ? 1 : 0;
for( size_t i = 0; i < matches.size(); i++ )
{
if( (int)matches[i].size() != knn )
badCount++;
else
{
int localBadCount = 0;
for( int k = 0; k < knn; k++ )
{
DMatch match = matches[i][k];
{
if( i < queryDescCount/2 )
{
if( (match.queryIdx != (int)i) || (match.trainIdx != (int)i*countFactor + k + shift) ||
(match.imgIdx != 0) )
localBadCount++;
}
else
{
if( (match.queryIdx != (int)i) || (match.trainIdx != ((int)i-queryDescCount/2)*countFactor + k + shift) ||
(match.imgIdx != 1) )
localBadCount++;
}
}
}
badCount += localBadCount > 0 ? 1 : 0;
}
}
if (badCount > 0)
{
ts->printf( cvtest::TS::LOG, "%f - too large bad matches part while test knnMatch() function (2).\n",
(float)badCount/(float)queryDescCount );
ts->set_failed_test_info( cvtest::TS::FAIL_BAD_ACCURACY );
}
}
}
}
void CV_GpuBruteForceMatcherTest::radiusMatchTest( const GpuMat& query, const GpuMat& train )
{
bool atomics_ok = TargetArchs::builtWith(GLOBAL_ATOMICS) && DeviceInfo().supports(GLOBAL_ATOMICS);
if (!atomics_ok)
{
ts->printf(cvtest::TS::CONSOLE, "\nCode and device atomics support is required for radiusMatch (CC >= 1.1)");
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
dmatcher.clear();
// test const version of match()
{
const float radius = 1.f/countFactor;
vector< vector<DMatch> > matches;
dmatcher.radiusMatch( query, train, matches, radius );
if( (int)matches.size() != queryDescCount )
{
ts->printf(cvtest::TS::LOG, "Incorrect matches count while test radiusMatch() function (1).\n");
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
else
{
int badCount = 0;
for( size_t i = 0; i < matches.size(); i++ )
{
if( (int)matches[i].size() != 1 )
badCount++;
else
{
DMatch match = matches[i][0];
if( (match.queryIdx != (int)i) || (match.trainIdx != (int)i*countFactor) || (match.imgIdx != 0) )
badCount++;
}
}
if (badCount > 0)
{
ts->printf( cvtest::TS::LOG, "%f - too large bad matches part while test radiusMatch() function (1).\n",
(float)badCount/(float)queryDescCount );
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
}
}
// test version of match() with add()
{
int n = 3;
const float radius = 1.f/countFactor * n;
vector< vector<DMatch> > matches;
// make add() twice to test such case
dmatcher.add( vector<GpuMat>(1,train.rowRange(0, train.rows/2)) );
dmatcher.add( vector<GpuMat>(1,train.rowRange(train.rows/2, train.rows)) );
// prepare masks (make first nearest match illegal)
vector<GpuMat> masks(2);
for(int mi = 0; mi < 2; mi++ )
{
masks[mi] = GpuMat(query.rows, train.rows/2, CV_8UC1, Scalar::all(1));
for( int di = 0; di < queryDescCount/2; di++ )
masks[mi].col(di*countFactor).setTo(Scalar::all(0));
}
dmatcher.radiusMatch( query, matches, radius, masks );
int curRes = cvtest::TS::OK;
if( (int)matches.size() != queryDescCount )
{
ts->printf(cvtest::TS::LOG, "Incorrect matches count while test radiusMatch() function (1).\n");
ts->set_failed_test_info( cvtest::TS::FAIL_INVALID_OUTPUT );
}
int badCount = 0;
int shift = dmatcher.isMaskSupported() ? 1 : 0;
int needMatchCount = dmatcher.isMaskSupported() ? n-1 : n;
for( size_t i = 0; i < matches.size(); i++ )
{
if( (int)matches[i].size() != needMatchCount )
badCount++;
else
{
int localBadCount = 0;
for( int k = 0; k < needMatchCount; k++ )
{
DMatch match = matches[i][k];
{
if( i < queryDescCount/2 )
{
if( (match.queryIdx != (int)i) || (match.trainIdx != (int)i*countFactor + k + shift) ||
(match.imgIdx != 0) )
localBadCount++;
}
else
{
if( (match.queryIdx != (int)i) || (match.trainIdx != ((int)i-queryDescCount/2)*countFactor + k + shift) ||
(match.imgIdx != 1) )
localBadCount++;
}
}
}
badCount += localBadCount > 0 ? 1 : 0;
}
}
if (badCount > 0)
{
curRes = cvtest::TS::FAIL_INVALID_OUTPUT;
ts->printf( cvtest::TS::LOG, "%f - too large bad matches part while test radiusMatch() function (2).\n",
(float)badCount/(float)queryDescCount );
ts->set_failed_test_info( cvtest::TS::FAIL_BAD_ACCURACY );
}
}
}
void CV_GpuBruteForceMatcherTest::dataTest(int dim)
{
GpuMat query, train;
generateData(query, train, dim);
matchTest(query, train);
knnMatchTest(query, train);
radiusMatchTest(query, train);
dmatcher.clear();
}
void CV_GpuBruteForceMatcherTest::run(int)
{
emptyDataTest();
dataTest(50);
dataTest(64);
dataTest(100);
dataTest(128);
dataTest(200);
dataTest(256);
dataTest(300);
}
TEST(BruteForceMatcher, accuracy) { CV_GpuBruteForceMatcherTest test; test.safe_run(); }

498
modules/gpu/test/test_calib3d.cpp
View File

@ -1,135 +1,363 @@
/*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.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
using namespace cv;
using namespace cv::gpu;
using namespace cvtest;
TEST(projectPoints, accuracy)
{
RNG& rng = TS::ptr()->get_rng();
Mat src = randomMat(rng, Size(1000, 1), CV_32FC3, 0, 10, false);
Mat rvec = randomMat(rng, Size(3, 1), CV_32F, 0, 1, false);
Mat tvec = randomMat(rng, Size(3, 1), CV_32F, 0, 1, false);
Mat camera_mat = randomMat(rng, Size(3, 3), CV_32F, 0, 1, false);
camera_mat.at<float>(0, 1) = 0.f;
camera_mat.at<float>(1, 0) = 0.f;
camera_mat.at<float>(2, 0) = 0.f;
camera_mat.at<float>(2, 1) = 0.f;
vector<Point2f> dst;
projectPoints(src, rvec, tvec, camera_mat, Mat(1, 8, CV_32F, Scalar::all(0)), dst);
GpuMat d_dst;
projectPoints(GpuMat(src), rvec, tvec, camera_mat, Mat(), d_dst);
ASSERT_EQ(dst.size(), (size_t)d_dst.cols);
ASSERT_EQ(1, d_dst.rows);
ASSERT_EQ(CV_32FC2, d_dst.type());
Mat h_dst(d_dst);
for (size_t i = 0; i < dst.size(); ++i)
{
Point2f res_gold = dst[i];
Point2f res_actual = h_dst.at<Point2f>(0, i);
Point2f err = res_actual - res_gold;
ASSERT_LT(err.dot(err) / res_gold.dot(res_gold), 1e-3f);
}
}
TEST(transformPoints, accuracy)
{
RNG& rng = TS::ptr()->get_rng();
Mat src = randomMat(rng, Size(1000, 1), CV_32FC3, 0, 10, false);
Mat rvec = randomMat(rng, Size(3, 1), CV_32F, 0, 1, false);
Mat tvec = randomMat(rng, Size(3, 1), CV_32F, 0, 1, false);
GpuMat d_dst;
transformPoints(GpuMat(src), rvec, tvec, d_dst);
ASSERT_TRUE(src.size() == d_dst.size());
ASSERT_EQ(src.type(), d_dst.type());
Mat h_dst(d_dst);
Mat rot;
Rodrigues(rvec, rot);
for (int i = 0; i < h_dst.cols; ++i)
{
Point3f p = src.at<Point3f>(0, i);
Point3f res_gold(
rot.at<float>(0, 0) * p.x + rot.at<float>(0, 1) * p.y + rot.at<float>(0, 2) * p.z + tvec.at<float>(0, 0),
rot.at<float>(1, 0) * p.x + rot.at<float>(1, 1) * p.y + rot.at<float>(1, 2) * p.z + tvec.at<float>(0, 1),
rot.at<float>(2, 0) * p.x + rot.at<float>(2, 1) * p.y + rot.at<float>(2, 2) * p.z + tvec.at<float>(0, 2));
Point3f res_actual = h_dst.at<Point3f>(0, i);
Point3f err = res_actual - res_gold;
ASSERT_LT(err.dot(err) / res_gold.dot(res_gold), 1e-3f);
}
}
TEST(solvePnPRansac, accuracy)
{
RNG& rng = TS::ptr()->get_rng();
const int num_points = 5000;
Mat object = randomMat(rng, Size(num_points, 1), CV_32FC3, 0, 100, false);
Mat camera_mat = randomMat(rng, Size(3, 3), CV_32F, 0.5, 1, false);
camera_mat.at<float>(0, 1) = 0.f;
camera_mat.at<float>(1, 0) = 0.f;
camera_mat.at<float>(2, 0) = 0.f;
camera_mat.at<float>(2, 1) = 0.f;
Mat rvec_gold = randomMat(rng, Size(3, 1), CV_32F, 0, 1, false);
Mat tvec_gold = randomMat(rng, Size(3, 1), CV_32F, 0, 1, false);
vector<Point2f> image_vec;
projectPoints(object, rvec_gold, tvec_gold, camera_mat, Mat(1, 8, CV_32F, Scalar::all(0)), image_vec);
Mat image(1, image_vec.size(), CV_32FC2, &image_vec[0]);
Mat rvec, tvec;
vector<int> inliers;
gpu::solvePnPRansac(object, image, camera_mat, Mat(1, 8, CV_32F, Scalar::all(0)), rvec, tvec, false, 200, 2.f, 100, &inliers);
ASSERT_LE(norm(rvec - rvec_gold), 1e-3f);
ASSERT_LE(norm(tvec - tvec_gold), 1e-3f);
}
/*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.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation may not be used to endorse or promote products
// derived from this software without specific prior written permission.
//
// This software is provided by the copyright holders and contributors "as is" and
// any express or implied warranties, including, but not limited to, the implied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "test_precomp.hpp"
#ifdef HAVE_CUDA
struct StereoTest : testing::TestWithParam<cv::gpu::DeviceInfo>
{
static cv::Mat img_l;
static cv::Mat img_r;
static cv::Mat img_template;
static void TearDownTestCase()
{
img_l.release();
img_r.release();
img_template.release();
}
cv::gpu::DeviceInfo devInfo;
virtual void SetUp()
{
devInfo = GetParam();
cv::gpu::setDevice(devInfo.deviceID());
}
};
cv::Mat StereoTest::img_l;
cv::Mat StereoTest::img_r;
cv::Mat StereoTest::img_template;
//////////////////////////////////////////////////////////////////////////
// BlockMatching
struct StereoBlockMatching : StereoTest
{
static void SetUpTestCase()
{
img_l = readImage("stereobm/aloe-L.png", CV_LOAD_IMAGE_GRAYSCALE);
img_r = readImage("stereobm/aloe-R.png", CV_LOAD_IMAGE_GRAYSCALE);
img_template = readImage("stereobm/aloe-disp.png", CV_LOAD_IMAGE_GRAYSCALE);
}
};
TEST_P(StereoBlockMatching, Regression)
{
ASSERT_TRUE(!img_l.empty() && !img_r.empty() && !img_template.empty());
PRINT_PARAM(devInfo);
cv::gpu::GpuMat disp;
cv::gpu::StereoBM_GPU bm(0, 128, 19);
bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
disp.convertTo(disp, img_template.type());
ASSERT_EQ(img_template.size(), disp.size());
double norm = cv::norm(img_template, (cv::Mat)disp, cv::NORM_INF);
ASSERT_EQ(0.0, norm);
}
INSTANTIATE_TEST_CASE_P(Calib3D, StereoBlockMatching, testing::ValuesIn(devices()));
//////////////////////////////////////////////////////////////////////////
// BeliefPropagation
struct StereoBeliefPropagation : StereoTest
{
static void SetUpTestCase()
{
img_l = readImage("stereobp/aloe-L.png");
img_r = readImage("stereobp/aloe-R.png");
img_template = readImage("stereobp/aloe-disp.png", CV_LOAD_IMAGE_GRAYSCALE);
}
};
TEST_P(StereoBeliefPropagation, Regression)
{
ASSERT_TRUE(!img_l.empty() && !img_r.empty() && !img_template.empty());
PRINT_PARAM(devInfo);
cv::gpu::GpuMat disp;
cv::gpu::StereoBeliefPropagation bpm(64, 8, 2, 25, 0.1f, 15, 1, CV_16S);
bpm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
disp.convertTo(disp, img_template.type());
ASSERT_EQ(img_template.size(), disp.size());
double norm = cv::norm(img_template, (cv::Mat)disp, cv::NORM_INF);
ASSERT_EQ(0.0, norm);
}
INSTANTIATE_TEST_CASE_P(Calib3D, StereoBeliefPropagation, testing::ValuesIn(devices()));
//////////////////////////////////////////////////////////////////////////
// ConstantSpaceBP
struct StereoConstantSpaceBP : StereoTest
{
static void SetUpTestCase()
{
img_l = readImage("csstereobp/aloe-L.png");
img_r = readImage("csstereobp/aloe-R.png");
}
virtual void SetUp()
{
StereoTest::SetUp();
if (supportFeature(GetParam(), cv::gpu::FEATURE_SET_COMPUTE_20))
img_template = readImage("csstereobp/aloe-disp.png", CV_LOAD_IMAGE_GRAYSCALE);
else
img_template = readImage("csstereobp/aloe-disp_CC1X.png", CV_LOAD_IMAGE_GRAYSCALE);
}
};
TEST_P(StereoConstantSpaceBP, Regression)
{
ASSERT_TRUE(!img_l.empty() && !img_r.empty() && !img_template.empty());
PRINT_PARAM(devInfo);
cv::gpu::GpuMat disp;
cv::gpu::StereoConstantSpaceBP bpm(128, 16, 4, 4);
bpm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
disp.convertTo(disp, img_template.type());
ASSERT_EQ(img_template.size(), disp.size());
double norm = cv::norm(img_template, (cv::Mat)disp, cv::NORM_INF);
ASSERT_EQ(0.0, norm);
}
INSTANTIATE_TEST_CASE_P(Calib3D, StereoConstantSpaceBP, testing::ValuesIn(devices()));
///////////////////////////////////////////////////////////////////////////////////////////////////////
// projectPoints
struct ProjectPoints : testing::TestWithParam<cv::gpu::DeviceInfo>
{
cv::gpu::DeviceInfo devInfo;
cv::Mat src;
cv::Mat rvec;
cv::Mat tvec;
cv::Mat camera_mat;
std::vector<cv::Point2f> dst_gold;
virtual void SetUp()
{
devInfo = GetParam();
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
src = cvtest::randomMat(rng, cv::Size(1000, 1), CV_32FC3, 0, 10, false);
rvec = cvtest::randomMat(rng, cv::Size(3, 1), CV_32F, 0, 1, false);
tvec = cvtest::randomMat(rng, cv::Size(3, 1), CV_32F, 0, 1, false);
camera_mat = cvtest::randomMat(rng, cv::Size(3, 3), CV_32F, 0, 1, false);
camera_mat.at<float>(0, 1) = 0.f;
camera_mat.at<float>(1, 0) = 0.f;
camera_mat.at<float>(2, 0) = 0.f;
camera_mat.at<float>(2, 1) = 0.f;
cv::projectPoints(src, rvec, tvec, camera_mat, cv::Mat(1, 8, CV_32F, cv::Scalar::all(0)), dst_gold);
}
};
TEST_P(ProjectPoints, Accuracy)
{
PRINT_PARAM(devInfo);
cv::Mat dst;
ASSERT_NO_THROW(
cv::gpu::GpuMat d_dst;
cv::gpu::projectPoints(cv::gpu::GpuMat(src), rvec, tvec, camera_mat, cv::Mat(), d_dst);
d_dst.download(dst);
);
ASSERT_EQ(dst_gold.size(), dst.cols);
ASSERT_EQ(1, dst.rows);
ASSERT_EQ(CV_32FC2, dst.type());
for (size_t i = 0; i < dst_gold.size(); ++i)
{
cv::Point2f res_gold = dst_gold[i];
cv::Point2f res_actual = dst.at<cv::Point2f>(0, i);
cv::Point2f err = res_actual - res_gold;
ASSERT_LE(err.dot(err) / res_gold.dot(res_gold), 1e-3f);
}
}
INSTANTIATE_TEST_CASE_P(Calib3D, ProjectPoints, testing::ValuesIn(devices()));
///////////////////////////////////////////////////////////////////////////////////////////////////////
// transformPoints
struct TransformPoints : testing::TestWithParam<cv::gpu::DeviceInfo>
{
cv::gpu::DeviceInfo devInfo;
cv::Mat src;
cv::Mat rvec;
cv::Mat tvec;
cv::Mat rot;
virtual void SetUp()
{
devInfo = GetParam();
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
src = cvtest::randomMat(rng, cv::Size(1000, 1), CV_32FC3, 0, 10, false);
rvec = cvtest::randomMat(rng, cv::Size(3, 1), CV_32F, 0, 1, false);
tvec = cvtest::randomMat(rng, cv::Size(3, 1), CV_32F, 0, 1, false);
cv::Rodrigues(rvec, rot);
}
};
TEST_P(TransformPoints, Accuracy)
{
PRINT_PARAM(devInfo);
cv::Mat dst;
ASSERT_NO_THROW(
cv::gpu::GpuMat d_dst;
cv::gpu::transformPoints(cv::gpu::GpuMat(src), rvec, tvec, d_dst);
d_dst.download(dst);
);
ASSERT_EQ(src.size(), dst.size());
ASSERT_EQ(src.type(), dst.type());
for (int i = 0; i < dst.cols; ++i)
{
cv::Point3f p = src.at<cv::Point3f>(0, i);
cv::Point3f res_gold(
rot.at<float>(0, 0) * p.x + rot.at<float>(0, 1) * p.y + rot.at<float>(0, 2) * p.z + tvec.at<float>(0, 0),
rot.at<float>(1, 0) * p.x + rot.at<float>(1, 1) * p.y + rot.at<float>(1, 2) * p.z + tvec.at<float>(0, 1),
rot.at<float>(2, 0) * p.x + rot.at<float>(2, 1) * p.y + rot.at<float>(2, 2) * p.z + tvec.at<float>(0, 2));
cv::Point3f res_actual = dst.at<cv::Point3f>(0, i);
cv::Point3f err = res_actual - res_gold;
ASSERT_LE(err.dot(err) / res_gold.dot(res_gold), 1e-3f);
}
}
INSTANTIATE_TEST_CASE_P(Calib3D, TransformPoints, testing::ValuesIn(devices()));
///////////////////////////////////////////////////////////////////////////////////////////////////////
// solvePnPRansac
struct SolvePnPRansac : testing::TestWithParam<cv::gpu::DeviceInfo>
{
static const int num_points = 5000;
cv::gpu::DeviceInfo devInfo;
cv::Mat object;
cv::Mat camera_mat;
std::vector<cv::Point2f> image_vec;
cv::Mat rvec_gold;
cv::Mat tvec_gold;
virtual void SetUp()
{
devInfo = GetParam();
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
object = cvtest::randomMat(rng, cv::Size(num_points, 1), CV_32FC3, 0, 100, false);
camera_mat = cvtest::randomMat(rng, cv::Size(3, 3), CV_32F, 0.5, 1, false);
camera_mat.at<float>(0, 1) = 0.f;
camera_mat.at<float>(1, 0) = 0.f;
camera_mat.at<float>(2, 0) = 0.f;
camera_mat.at<float>(2, 1) = 0.f;
rvec_gold = cvtest::randomMat(rng, cv::Size(3, 1), CV_32F, 0, 1, false);
tvec_gold = cvtest::randomMat(rng, cv::Size(3, 1), CV_32F, 0, 1, false);
cv::projectPoints(object, rvec_gold, tvec_gold, camera_mat, cv::Mat(1, 8, CV_32F, cv::Scalar::all(0)), image_vec);
}
};
TEST_P(SolvePnPRansac, Accuracy)
{
PRINT_PARAM(devInfo);
cv::Mat rvec, tvec;
std::vector<int> inliers;
ASSERT_NO_THROW(
cv::gpu::solvePnPRansac(object, cv::Mat(1, image_vec.size(), CV_32FC2, &image_vec[0]), camera_mat,
cv::Mat(1, 8, CV_32F, cv::Scalar::all(0)), rvec, tvec, false, 200, 2.f, 100, &inliers);
);
ASSERT_LE(cv::norm(rvec - rvec_gold), 1e-3f);
ASSERT_LE(cv::norm(tvec - tvec_gold), 1e-3f);
}
INSTANTIATE_TEST_CASE_P(Calib3D, SolvePnPRansac, testing::ValuesIn(devices()));
#endif // HAVE_CUDA

399
modules/gpu/test/test_dft_routines.cpp
View File

@ -1,399 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
using namespace cv;
using namespace cv::gpu;
using namespace std;
struct CV_GpuMulSpectrumsTest: cvtest::BaseTest
{
CV_GpuMulSpectrumsTest() {}
void run(int)
{
test(0);
testConj(0);
testScaled(0);
testScaledConj(0);
test(DFT_ROWS);
testConj(DFT_ROWS);
testScaled(DFT_ROWS);
testScaledConj(DFT_ROWS);
}
void gen(int cols, int rows, Mat& mat)
{
RNG rng;
mat.create(rows, cols, CV_32FC2);
rng.fill(mat, RNG::UNIFORM, Scalar::all(0.f), Scalar::all(10.f));
}
bool cmp(const Mat& gold, const Mat& mine, float max_err=1e-3f)
{
if (gold.size() != mine.size())
{
ts->printf(cvtest::TS::CONSOLE, "bad sizes: gold: %d d%, mine: %d %d\n", gold.cols, gold.rows, mine.cols, mine.rows);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
if (gold.type() != mine.type())
{
ts->printf(cvtest::TS::CONSOLE, "bad types: gold=%d, mine=%d\n", gold.type(), mine.type());
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
for (int i = 0; i < gold.rows; ++i)
{
for (int j = 0; j < gold.cols * 2; ++j)
{
float gold_ = gold.at<float>(i, j);
float mine_ = mine.at<float>(i, j);
if (fabs(gold_ - mine_) > max_err)
{
ts->printf(cvtest::TS::CONSOLE, "bad values at %d %d: gold=%f, mine=%f\n", j, i, gold_, mine_);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
}
}
return true;
}
bool cmpScaled(const Mat& gold, const Mat& mine, float scale, float max_err=1e-3f)
{
if (gold.size() != mine.size())
{
ts->printf(cvtest::TS::CONSOLE, "bad sizes: gold: %d d%, mine: %d %d\n", gold.cols, gold.rows, mine.cols, mine.rows);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
if (gold.type() != mine.type())
{
ts->printf(cvtest::TS::CONSOLE, "bad types: gold=%d, mine=%d\n", gold.type(), mine.type());
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
for (int i = 0; i < gold.rows; ++i)
{
for (int j = 0; j < gold.cols * 2; ++j)
{
float gold_ = gold.at<float>(i, j) * scale;
float mine_ = mine.at<float>(i, j);
if (fabs(gold_ - mine_) > max_err)
{
ts->printf(cvtest::TS::CONSOLE, "bad values at %d %d: gold=%f, mine=%f\n", j, i, gold_, mine_);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
}
}
return true;
}
void test(int flags)
{
int cols = 1 + rand() % 100, rows = 1 + rand() % 1000;
Mat a, b;
gen(cols, rows, a);
gen(cols, rows, b);
Mat c_gold;
mulSpectrums(a, b, c_gold, flags, false);
GpuMat d_c;
mulSpectrums(GpuMat(a), GpuMat(b), d_c, flags, false);
if (!cmp(c_gold, Mat(d_c)))
ts->printf(cvtest::TS::CONSOLE, "test failed: cols=%d, rows=%d, flags=%d\n", cols, rows, flags);
}
void testConj(int flags)
{
int cols = 1 + rand() % 100, rows = 1 + rand() % 1000;
Mat a, b;
gen(cols, rows, a);
gen(cols, rows, b);
Mat c_gold;
mulSpectrums(a, b, c_gold, flags, true);
GpuMat d_c;
mulSpectrums(GpuMat(a), GpuMat(b), d_c, flags, true);
if (!cmp(c_gold, Mat(d_c)))
ts->printf(cvtest::TS::CONSOLE, "testConj failed: cols=%d, rows=%d, flags=%d\n", cols, rows, flags);
}
void testScaled(int flags)
{
int cols = 1 + rand() % 100, rows = 1 + rand() % 1000;
Mat a, b;
gen(cols, rows, a);
gen(cols, rows, b);
float scale = 1.f / a.size().area();
Mat c_gold;
mulSpectrums(a, b, c_gold, flags, false);
GpuMat d_c;
mulAndScaleSpectrums(GpuMat(a), GpuMat(b), d_c, flags, scale, false);
if (!cmpScaled(c_gold, Mat(d_c), scale))
ts->printf(cvtest::TS::CONSOLE, "testScaled failed: cols=%d, rows=%d, flags=%d\n", cols, rows, flags);
}
void testScaledConj(int flags)
{
int cols = 1 + rand() % 100, rows = 1 + rand() % 1000;
Mat a, b;
gen(cols, rows, a);
gen(cols, rows, b);
float scale = 1.f / a.size().area();
Mat c_gold;
mulSpectrums(a, b, c_gold, flags, true);
GpuMat d_c;
mulAndScaleSpectrums(GpuMat(a), GpuMat(b), d_c, flags, scale, true);
if (!cmpScaled(c_gold, Mat(d_c), scale))
ts->printf(cvtest::TS::CONSOLE, "testScaledConj failed: cols=%d, rows=%d, flags=%D\n", cols, rows, flags);
}
} CV_GpuMulSpectrumsTest_inst;
struct CV_GpuDftTest: cvtest::BaseTest
{
CV_GpuDftTest() {}
void run(int)
{
srand(0);
int cols = 2 + rand() % 100, rows = 2 + rand() % 100;
for (int i = 0; i < 2; ++i)
{
bool inplace = i != 0;
testC2C("no flags", cols, rows, 0, inplace);
testC2C("no flags 0 1", cols, rows + 1, 0, inplace);
testC2C("no flags 1 0", cols, rows + 1, 0, inplace);
testC2C("no flags 1 1", cols + 1, rows, 0, inplace);
testC2C("DFT_INVERSE", cols, rows, DFT_INVERSE, inplace);
testC2C("DFT_ROWS", cols, rows, DFT_ROWS, inplace);
testC2C("single col", 1, rows, 0, inplace);
testC2C("single row", cols, 1, 0, inplace);
testC2C("single col inversed", 1, rows, DFT_INVERSE, inplace);
testC2C("single row inversed", cols, 1, DFT_INVERSE, inplace);
testC2C("single row DFT_ROWS", cols, 1, DFT_ROWS, inplace);
testC2C("size 1 2", 1, 2, 0, inplace);
testC2C("size 2 1", 2, 1, 0, inplace);
}
testR2CThenC2R("sanity", cols, rows);
testR2CThenC2R("sanity 0 1", cols, rows + 1);
testR2CThenC2R("sanity 1 0", cols + 1, rows);
testR2CThenC2R("sanity 1 1", cols + 1, rows + 1);
testR2CThenC2R("single col", 1, rows);
testR2CThenC2R("single col 1", 1, rows + 1);
testR2CThenC2R("single row", cols, 1);
testR2CThenC2R("single row 1", cols + 1, 1);
testR2CThenC2R("sanity", cols, rows, true);
testR2CThenC2R("sanity 0 1", cols, rows + 1, true);
testR2CThenC2R("sanity 1 0", cols + 1, rows, true);
testR2CThenC2R("sanity 1 1", cols + 1, rows + 1, true);
testR2CThenC2R("single row", cols, 1, true);
testR2CThenC2R("single row 1", cols + 1, 1, true);
}
void gen(int cols, int rows, int cn, Mat& mat)
{
RNG rng(1);
mat.create(rows, cols, CV_MAKETYPE(CV_32F, cn));
rng.fill(mat, RNG::UNIFORM, Scalar::all(0.f), Scalar::all(10.f));
}
bool cmp(const Mat& gold, const Mat& mine, float max_err=1e-3f)
{
if (gold.size() != mine.size())
{
ts->printf(cvtest::TS::CONSOLE, "bad sizes: gold: %d %d, mine: %d %d\n", gold.cols, gold.rows, mine.cols, mine.rows);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
if (gold.depth() != mine.depth())
{
ts->printf(cvtest::TS::CONSOLE, "bad depth: gold=%d, mine=%d\n", gold.depth(), mine.depth());
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
if (gold.channels() != mine.channels())
{
ts->printf(cvtest::TS::CONSOLE, "bad channel count: gold=%d, mine=%d\n", gold.channels(), mine.channels());
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
for (int i = 0; i < gold.rows; ++i)
{
for (int j = 0; j < gold.cols * gold.channels(); ++j)
{
float gold_ = gold.at<float>(i, j);
float mine_ = mine.at<float>(i, j);
if (fabs(gold_ - mine_) > max_err)
{
ts->printf(cvtest::TS::CONSOLE, "bad values at %d %d: gold=%f, mine=%f\n", j / gold.channels(), i, gold_, mine_);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
}
}
return true;
}
void testC2C(const std::string& hint, int cols, int rows, int flags, bool inplace=false)
{
Mat a;
gen(cols, rows, 2, a);
Mat b_gold;
dft(a, b_gold, flags);
GpuMat d_b;
GpuMat d_b_data;
if (inplace)
{
d_b_data.create(1, a.size().area(), CV_32FC2);
d_b = GpuMat(a.rows, a.cols, CV_32FC2, d_b_data.ptr(), a.cols * d_b_data.elemSize());
}
dft(GpuMat(a), d_b, Size(cols, rows), flags);
bool ok = true;
if (ok && inplace && d_b.ptr() != d_b_data.ptr())
{
ts->printf(cvtest::TS::CONSOLE, "unnecessary reallocation was done\n");
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
ok = false;
}
if (ok && d_b.depth() != CV_32F)
{
ts->printf(cvtest::TS::CONSOLE, "bad depth: %d\n", d_b.depth());
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
ok = false;
}
if (ok && d_b.channels() != 2)
{
ts->printf(cvtest::TS::CONSOLE, "bad channel count: %d\n", d_b.channels());
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
ok = false;
}
if (ok) ok = cmp(b_gold, Mat(d_b), rows * cols * 1e-4f);
if (!ok)
ts->printf(cvtest::TS::CONSOLE, "testC2C failed: hint=%s, cols=%d, rows=%d, flags=%d, inplace=%d\n",
hint.c_str(), cols, rows, flags, inplace);
}
void testR2CThenC2R(const std::string& hint, int cols, int rows, bool inplace=false)
{
Mat a;
gen(cols, rows, 1, a);
bool ok = true;
GpuMat d_b, d_c;
GpuMat d_b_data, d_c_data;
if (inplace)
{
if (a.cols == 1)
{
d_b_data.create(1, (a.rows / 2 + 1) * a.cols, CV_32FC2);
d_b = GpuMat(a.rows / 2 + 1, a.cols, CV_32FC2, d_b_data.ptr(), a.cols * d_b_data.elemSize());
}
else
{
d_b_data.create(1, a.rows * (a.cols / 2 + 1), CV_32FC2);
d_b = GpuMat(a.rows, a.cols / 2 + 1, CV_32FC2, d_b_data.ptr(), (a.cols / 2 + 1) * d_b_data.elemSize());
}
d_c_data.create(1, a.size().area(), CV_32F);
d_c = GpuMat(a.rows, a.cols, CV_32F, d_c_data.ptr(), a.cols * d_c_data.elemSize());
}
dft(GpuMat(a), d_b, Size(cols, rows), 0);
dft(d_b, d_c, Size(cols, rows), DFT_REAL_OUTPUT | DFT_SCALE);
if (ok && inplace && d_b.ptr() != d_b_data.ptr())
{
ts->printf(cvtest::TS::CONSOLE, "unnecessary reallocation was done for b\n");
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
ok = false;
}
if (ok && inplace && d_c.ptr() != d_c_data.ptr())
{
ts->printf(cvtest::TS::CONSOLE, "unnecessary reallocation was done for c\n");
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
ok = false;
}
if (ok && d_c.depth() != CV_32F)
{
ts->printf(cvtest::TS::CONSOLE, "bad depth: %d\n", d_c.depth());
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
ok = false;
}
if (ok && d_c.channels() != 1)
{
ts->printf(cvtest::TS::CONSOLE, "bad channel count: %d\n", d_c.channels());
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
ok = false;
}
if (ok) ok = cmp(a, Mat(d_c), rows * cols * 1e-5f);
if (!ok)
ts->printf(cvtest::TS::CONSOLE, "testR2CThenC2R failed: hint=%s, cols=%d, rows=%d, inplace=%d\n",
hint.c_str(), cols, rows, inplace);
}
};
TEST(dft, accuracy) { CV_GpuDftTest test; test.safe_run(); }

1389
modules/gpu/test/test_features2d.cpp
View File

File diff suppressed because it is too large Load Diff

675
modules/gpu/test/test_filters.cpp
View File

@ -39,329 +39,516 @@
//
//M*/
#include <iostream>
#include <cmath>
#include <limits>
#include "test_precomp.hpp"
using namespace cv;
using namespace std;
using namespace gpu;
#ifdef HAVE_CUDA
class CV_GpuNppFilterTest : public cvtest::BaseTest
struct FilterTest
{
public:
CV_GpuNppFilterTest(const char* /*test_name*/, const char* /*test_funcs*/) {}
virtual ~CV_GpuNppFilterTest() {}
static cv::Mat img_rgba;
static cv::Mat img_gray;
protected:
void run(int);
virtual int test(const Mat& img) = 0;
int test8UC1(const Mat& img)
static void SetUpTestCase()
{
cv::Mat img_C1;
cvtColor(img, img_C1, CV_BGR2GRAY);
return test(img_C1);
cv::Mat img = readImage("stereobp/aloe-L.png");
cv::cvtColor(img, img_rgba, CV_BGR2BGRA);
cv::cvtColor(img, img_gray, CV_BGR2GRAY);
}
int test8UC4(const Mat& img)
static void TearDownTestCase()
{
cv::Mat img_C4;
cvtColor(img, img_C4, CV_BGR2BGRA);
return test(img_C4);
}
int CheckNorm(const Mat& m1, const Mat& m2, const Size& ksize)
{
Rect roi = Rect(ksize.width, ksize.height, m1.cols - 2 * ksize.width, m1.rows - 2 * ksize.height);
Mat m1ROI = m1(roi);
Mat m2ROI = m2(roi);
double res = norm(m1ROI, m2ROI, NORM_INF);
// Max difference (2.0) in GaussianBlur
if (res <= 2)
return cvtest::TS::OK;
ts->printf(cvtest::TS::LOG, "Norm: %f\n", res);
return cvtest::TS::FAIL_GENERIC;
img_rgba.release();
img_gray.release();
}
};
void CV_GpuNppFilterTest::run( int )
{
cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");
if (img.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
cv::Mat FilterTest::img_rgba;
cv::Mat FilterTest::img_gray;
//run tests
int testResult = cvtest::TS::OK;
if (test8UC1(img) != cvtest::TS::OK)
testResult = cvtest::TS::FAIL_GENERIC;
if (test8UC4(img) != cvtest::TS::OK)
testResult = cvtest::TS::FAIL_GENERIC;
ts->set_failed_test_info(testResult);
ts->set_failed_test_info(cvtest::TS::OK);
static double checkNorm(const cv::Mat& m1, const cv::Mat& m2, const cv::Size& ksize)
{
cv::Rect roi(ksize.width, ksize.height, m1.cols - 2 * ksize.width, m1.rows - 2 * ksize.height);
cv::Mat m1ROI = m1(roi);
cv::Mat m2ROI = m2(roi);
return checkNorm(m1ROI, m2ROI);
}
////////////////////////////////////////////////////////////////////////////////
static double checkNorm(const cv::Mat& m1, const cv::Mat& m2, int ksize)
{
return checkNorm(m1, m2, cv::Size(ksize, ksize));
}
#define EXPECT_MAT_NEAR_KSIZE(mat1, mat2, ksize, eps) \
{ \
ASSERT_EQ(mat1.type(), mat2.type()); \
ASSERT_EQ(mat1.size(), mat2.size()); \
EXPECT_LE(checkNorm(mat1, mat2, ksize), eps); \
}
/////////////////////////////////////////////////////////////////////////////////////////////////
// blur
struct CV_GpuNppImageBlurTest : public CV_GpuNppFilterTest
struct Blur : FilterTest, testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int, int> >
{
CV_GpuNppImageBlurTest() : CV_GpuNppFilterTest( "GPU-NppImageBlur", "blur" ) {}
cv::gpu::DeviceInfo devInfo;
cv::Size ksize;
int test(const Mat& img)
cv::Mat dst_gold_rgba;
cv::Mat dst_gold_gray;
using FilterTest::SetUpTestCase;
using FilterTest::TearDownTestCase;
virtual void SetUp()
{
int ksizes[] = {3, 5, 7};
int ksizes_num = sizeof(ksizes) / sizeof(int);
devInfo = std::tr1::get<0>(GetParam());
ksize = cv::Size(std::tr1::get<1>(GetParam()), std::tr1::get<2>(GetParam()));
int test_res = cvtest::TS::OK;
cv::gpu::setDevice(devInfo.deviceID());
for (int i = 0; i < ksizes_num; ++i)
{
for (int j = 0; j < ksizes_num; ++j)
{
Size ksize(ksizes[i], ksizes[j]);
ts->printf(cvtest::TS::LOG, "\nksize = (%dx%d)\n", ksizes[i], ksizes[j]);
Mat cpudst;
cv::blur(img, cpudst, ksize);
GpuMat gpu1(img);
GpuMat gpudst;
cv::gpu::blur(gpu1, gpudst, ksize);
if (CheckNorm(cpudst, gpudst, ksize) != cvtest::TS::OK)
test_res = cvtest::TS::FAIL_GENERIC;
}
}
return test_res;
cv::blur(img_rgba, dst_gold_rgba, ksize);
cv::blur(img_gray, dst_gold_gray, ksize);
}
};
////////////////////////////////////////////////////////////////////////////////
// Sobel
struct CV_GpuNppImageSobelTest : public CV_GpuNppFilterTest
TEST_P(Blur, Accuracy)
{
CV_GpuNppImageSobelTest() : CV_GpuNppFilterTest( "GPU-NppImageSobel", "Sobel" ) {}
ASSERT_TRUE(!img_rgba.empty() && !img_gray.empty());
int test(const Mat& img)
PRINT_PARAM(devInfo);
PRINT_PARAM(ksize);
cv::Mat dst_rgba;
cv::Mat dst_gray;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_dst_rgba;
cv::gpu::GpuMat dev_dst_gray;
cv::gpu::blur(cv::gpu::GpuMat(img_rgba), dev_dst_rgba, ksize);
cv::gpu::blur(cv::gpu::GpuMat(img_gray), dev_dst_gray, ksize);
dev_dst_rgba.download(dst_rgba);
dev_dst_gray.download(dst_gray);
);
EXPECT_MAT_NEAR_KSIZE(dst_gold_rgba, dst_rgba, ksize, 1.0);
EXPECT_MAT_NEAR_KSIZE(dst_gold_gray, dst_gray, ksize, 1.0);
}
INSTANTIATE_TEST_CASE_P(Filter, Blur, testing::Combine(
testing::ValuesIn(devices()),
testing::Values(3, 5, 7),
testing::Values(3, 5, 7)));
/////////////////////////////////////////////////////////////////////////////////////////////////
// sobel
struct Sobel : FilterTest, testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int, std::pair<int, int> > >
{
cv::gpu::DeviceInfo devInfo;
int ksize;
int dx, dy;
cv::Mat dst_gold_rgba;
cv::Mat dst_gold_gray;
using FilterTest::SetUpTestCase;
using FilterTest::TearDownTestCase;
virtual void SetUp()
{
int ksizes[] = {3, 5, 7};
int ksizes_num = sizeof(ksizes) / sizeof(int);
devInfo = std::tr1::get<0>(GetParam());
ksize = std::tr1::get<1>(GetParam());
std::pair<int, int> d = std::tr1::get<2>(GetParam());
dx = d.first; dy = d.second;
int dx = 1, dy = 0;
int test_res = cvtest::TS::OK;
for (int i = 0; i < ksizes_num; ++i)
{
ts->printf(cvtest::TS::LOG, "\nksize = %d\n", ksizes[i]);
Mat cpudst;
cv::Sobel(img, cpudst, -1, dx, dy, ksizes[i]);
GpuMat gpu1(img);
GpuMat gpudst;
cv::gpu::Sobel(gpu1, gpudst, -1, dx, dy, ksizes[i]);
if (CheckNorm(cpudst, gpudst, Size(ksizes[i], ksizes[i])) != cvtest::TS::OK)
test_res = cvtest::TS::FAIL_GENERIC;
}
return test_res;
cv::gpu::setDevice(devInfo.deviceID());
cv::Sobel(img_rgba, dst_gold_rgba, -1, dx, dy, ksize);
cv::Sobel(img_gray, dst_gold_gray, -1, dx, dy, ksize);
}
};
////////////////////////////////////////////////////////////////////////////////
// Scharr
struct CV_GpuNppImageScharrTest : public CV_GpuNppFilterTest
TEST_P(Sobel, Accuracy)
{
CV_GpuNppImageScharrTest() : CV_GpuNppFilterTest( "GPU-NppImageScharr", "Scharr" ) {}
ASSERT_TRUE(!img_rgba.empty() && !img_gray.empty());
int test(const Mat& img)
PRINT_PARAM(devInfo);
PRINT_PARAM(ksize);
PRINT_PARAM(dx);
PRINT_PARAM(dy);
cv::Mat dst_rgba;
cv::Mat dst_gray;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_dst_rgba;
cv::gpu::GpuMat dev_dst_gray;
cv::gpu::Sobel(cv::gpu::GpuMat(img_rgba), dev_dst_rgba, -1, dx, dy, ksize);
cv::gpu::Sobel(cv::gpu::GpuMat(img_gray), dev_dst_gray, -1, dx, dy, ksize);
dev_dst_rgba.download(dst_rgba);
dev_dst_gray.download(dst_gray);
);
EXPECT_MAT_NEAR_KSIZE(dst_gold_rgba, dst_rgba, ksize, 0.0);
EXPECT_MAT_NEAR_KSIZE(dst_gold_gray, dst_gray, ksize, 0.0);
}
INSTANTIATE_TEST_CASE_P(Filter, Sobel, testing::Combine(
testing::ValuesIn(devices()),
testing::Values(3, 5, 7),
testing::Values(std::make_pair(1, 0), std::make_pair(0, 1), std::make_pair(1, 1), std::make_pair(2, 0), std::make_pair(2, 1), std::make_pair(0, 2), std::make_pair(1, 2), std::make_pair(2, 2))));
/////////////////////////////////////////////////////////////////////////////////////////////////
// scharr
struct Scharr : FilterTest, testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, std::pair<int, int> > >
{
cv::gpu::DeviceInfo devInfo;
int dx, dy;
cv::Mat dst_gold_rgba;
cv::Mat dst_gold_gray;
using FilterTest::SetUpTestCase;
using FilterTest::TearDownTestCase;
virtual void SetUp()
{
int dx = 1, dy = 0;
devInfo = std::tr1::get<0>(GetParam());
std::pair<int, int> d = std::tr1::get<1>(GetParam());
dx = d.first; dy = d.second;
Mat cpudst;
cv::Scharr(img, cpudst, -1, dx, dy);
cv::gpu::setDevice(devInfo.deviceID());
GpuMat gpu1(img);
GpuMat gpudst;
cv::gpu::Scharr(gpu1, gpudst, -1, dx, dy);
return CheckNorm(cpudst, gpudst, Size(3, 3));
cv::Scharr(img_rgba, dst_gold_rgba, -1, dx, dy);
cv::Scharr(img_gray, dst_gold_gray, -1, dx, dy);
}
};
////////////////////////////////////////////////////////////////////////////////
// GaussianBlur
struct CV_GpuNppImageGaussianBlurTest : public CV_GpuNppFilterTest
TEST_P(Scharr, Accuracy)
{
CV_GpuNppImageGaussianBlurTest() : CV_GpuNppFilterTest( "GPU-NppImageGaussianBlur", "GaussianBlur" ) {}
ASSERT_TRUE(!img_rgba.empty() && !img_gray.empty());
int test(const Mat& img)
PRINT_PARAM(devInfo);
PRINT_PARAM(dx);
PRINT_PARAM(dy);
cv::Mat dst_rgba;
cv::Mat dst_gray;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_dst_rgba;
cv::gpu::GpuMat dev_dst_gray;
cv::gpu::Scharr(cv::gpu::GpuMat(img_rgba), dev_dst_rgba, -1, dx, dy);
cv::gpu::Scharr(cv::gpu::GpuMat(img_gray), dev_dst_gray, -1, dx, dy);
dev_dst_rgba.download(dst_rgba);
dev_dst_gray.download(dst_gray);
);
EXPECT_MAT_NEAR_KSIZE(dst_gold_rgba, dst_rgba, 3, 0.0);
EXPECT_MAT_NEAR_KSIZE(dst_gold_gray, dst_gray, 3, 0.0);
}
INSTANTIATE_TEST_CASE_P(Filter, Scharr, testing::Combine(
testing::ValuesIn(devices()),
testing::Values(std::make_pair(1, 0), std::make_pair(0, 1))));
/////////////////////////////////////////////////////////////////////////////////////////////////
// gaussianBlur
struct GaussianBlur : FilterTest, testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int, int> >
{
cv::gpu::DeviceInfo devInfo;
cv::Size ksize;
double sigma1, sigma2;
cv::Mat dst_gold_rgba;
cv::Mat dst_gold_gray;
using FilterTest::SetUpTestCase;
using FilterTest::TearDownTestCase;
virtual void SetUp()
{
int ksizes[] = {3, 5, 7};
int ksizes_num = sizeof(ksizes) / sizeof(int);
devInfo = std::tr1::get<0>(GetParam());
ksize = cv::Size(std::tr1::get<1>(GetParam()), std::tr1::get<2>(GetParam()));
int test_res = cvtest::TS::OK;
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
const double sigma1 = 3.0;
for (int i = 0; i < ksizes_num; ++i)
{
for (int j = 0; j < ksizes_num; ++j)
{
cv::Size ksize(ksizes[i], ksizes[j]);
ts->printf(cvtest::TS::LOG, "ksize = (%dx%d)\t\n", ksizes[i], ksizes[j]);
Mat cpudst;
cv::GaussianBlur(img, cpudst, ksize, sigma1);
GpuMat gpu1(img);
GpuMat gpudst;
cv::gpu::GaussianBlur(gpu1, gpudst, ksize, sigma1);
if (CheckNorm(cpudst, gpudst, ksize) != cvtest::TS::OK)
test_res = cvtest::TS::FAIL_GENERIC;
}
}
return test_res;
sigma1 = rng.uniform(0.1, 1.0);
sigma2 = rng.uniform(0.1, 1.0);
cv::GaussianBlur(img_rgba, dst_gold_rgba, ksize, sigma1, sigma2);
cv::GaussianBlur(img_gray, dst_gold_gray, ksize, sigma1, sigma2);
}
};
////////////////////////////////////////////////////////////////////////////////
// Laplacian
struct CV_GpuNppImageLaplacianTest : public CV_GpuNppFilterTest
TEST_P(GaussianBlur, Accuracy)
{
CV_GpuNppImageLaplacianTest() : CV_GpuNppFilterTest( "GPU-NppImageLaplacian", "Laplacian" ) {}
ASSERT_TRUE(!img_rgba.empty() && !img_gray.empty());
int test(const Mat& img)
PRINT_PARAM(devInfo);
PRINT_PARAM(ksize);
PRINT_PARAM(sigma1);
PRINT_PARAM(sigma2);
cv::Mat dst_rgba;
cv::Mat dst_gray;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_dst_rgba;
cv::gpu::GpuMat dev_dst_gray;
cv::gpu::GaussianBlur(cv::gpu::GpuMat(img_rgba), dev_dst_rgba, ksize, sigma1, sigma2);
cv::gpu::GaussianBlur(cv::gpu::GpuMat(img_gray), dev_dst_gray, ksize, sigma1, sigma2);
dev_dst_rgba.download(dst_rgba);
dev_dst_gray.download(dst_gray);
);
EXPECT_MAT_NEAR_KSIZE(dst_gold_rgba, dst_rgba, ksize, 3.0);
EXPECT_MAT_NEAR_KSIZE(dst_gold_gray, dst_gray, ksize, 3.0);
}
INSTANTIATE_TEST_CASE_P(Filter, GaussianBlur, testing::Combine(
testing::ValuesIn(devices()),
testing::Values(3, 5, 7),
testing::Values(3, 5, 7)));
/////////////////////////////////////////////////////////////////////////////////////////////////
// laplacian
struct Laplacian : FilterTest, testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int> >
{
cv::gpu::DeviceInfo devInfo;
int ksize;
cv::Mat dst_gold_rgba;
cv::Mat dst_gold_gray;
using FilterTest::SetUpTestCase;
using FilterTest::TearDownTestCase;
virtual void SetUp()
{
int ksizes[] = {1, 3};
int ksizes_num = sizeof(ksizes) / sizeof(int);
devInfo = std::tr1::get<0>(GetParam());
ksize = std::tr1::get<1>(GetParam());
int test_res = cvtest::TS::OK;
cv::gpu::setDevice(devInfo.deviceID());
for (int i = 0; i < ksizes_num; ++i)
{
ts->printf(cvtest::TS::LOG, "\nksize = %d\n", ksizes[i]);
Mat cpudst;
cv::Laplacian(img, cpudst, -1, ksizes[i]);
GpuMat gpu1(img);
GpuMat gpudst;
cv::gpu::Laplacian(gpu1, gpudst, -1, ksizes[i]);
if (CheckNorm(cpudst, gpudst, Size(3, 3)) != cvtest::TS::OK)
test_res = cvtest::TS::FAIL_GENERIC;
}
return test_res;
cv::Laplacian(img_rgba, dst_gold_rgba, -1, ksize);
cv::Laplacian(img_gray, dst_gold_gray, -1, ksize);
}
};
////////////////////////////////////////////////////////////////////////////////
// Erode
class CV_GpuErodeTest : public CV_GpuNppFilterTest
TEST_P(Laplacian, Accuracy)
{
public:
CV_GpuErodeTest() : CV_GpuNppFilterTest( "GPU-NppErode", "erode" ) {}
ASSERT_TRUE(!img_rgba.empty() && !img_gray.empty());
protected:
virtual int test(const Mat& img)
PRINT_PARAM(devInfo);
PRINT_PARAM(ksize);
cv::Mat dst_rgba;
cv::Mat dst_gray;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_dst_rgba;
cv::gpu::GpuMat dev_dst_gray;
cv::gpu::Laplacian(cv::gpu::GpuMat(img_rgba), dev_dst_rgba, -1, ksize);
cv::gpu::Laplacian(cv::gpu::GpuMat(img_gray), dev_dst_gray, -1, ksize);
dev_dst_rgba.download(dst_rgba);
dev_dst_gray.download(dst_gray);
);
EXPECT_MAT_NEAR_KSIZE(dst_gold_rgba, dst_rgba, 3, 0.0);
EXPECT_MAT_NEAR_KSIZE(dst_gold_gray, dst_gray, 3, 0.0);
}
INSTANTIATE_TEST_CASE_P(Filter, Laplacian, testing::Combine(
testing::ValuesIn(devices()),
testing::Values(1, 3)));
/////////////////////////////////////////////////////////////////////////////////////////////////
// erode
struct Erode : FilterTest, testing::TestWithParam<cv::gpu::DeviceInfo>
{
cv::gpu::DeviceInfo devInfo;
cv::Mat kernel;
cv::Mat dst_gold_rgba;
cv::Mat dst_gold_gray;
using FilterTest::SetUpTestCase;
using FilterTest::TearDownTestCase;
virtual void SetUp()
{
Mat kernel(Mat::ones(3, 3, CV_8U));
devInfo = GetParam();
cv::Mat cpuRes;
cv::erode(img, cpuRes, kernel);
cv::gpu::setDevice(devInfo.deviceID());
GpuMat gpuRes;
cv::gpu::erode(GpuMat(img), gpuRes, kernel);
kernel = cv::Mat::ones(3, 3, CV_8U);
return CheckNorm(cpuRes, gpuRes, Size(3, 3));
cv::erode(img_rgba, dst_gold_rgba, kernel);
cv::erode(img_gray, dst_gold_gray, kernel);
}
};
////////////////////////////////////////////////////////////////////////////////
// Dilate
class CV_GpuDilateTest : public CV_GpuNppFilterTest
TEST_P(Erode, Accuracy)
{
public:
CV_GpuDilateTest() : CV_GpuNppFilterTest( "GPU-NppDilate", "dilate" ) {}
ASSERT_TRUE(!img_rgba.empty() && !img_gray.empty());
protected:
virtual int test(const Mat& img)
PRINT_PARAM(devInfo);
cv::Mat dst_rgba;
cv::Mat dst_gray;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_dst_rgba;
cv::gpu::GpuMat dev_dst_gray;
cv::gpu::erode(cv::gpu::GpuMat(img_rgba), dev_dst_rgba, kernel);
cv::gpu::erode(cv::gpu::GpuMat(img_gray), dev_dst_gray, kernel);
dev_dst_rgba.download(dst_rgba);
dev_dst_gray.download(dst_gray);
);
EXPECT_MAT_NEAR_KSIZE(dst_gold_rgba, dst_rgba, 3, 0.0);
EXPECT_MAT_NEAR_KSIZE(dst_gold_gray, dst_gray, 3, 0.0);
}
INSTANTIATE_TEST_CASE_P(Filter, Erode, testing::ValuesIn(devices()));
/////////////////////////////////////////////////////////////////////////////////////////////////
// dilate
struct Dilate : FilterTest, testing::TestWithParam<cv::gpu::DeviceInfo>
{
cv::gpu::DeviceInfo devInfo;
cv::Mat kernel;
cv::Mat dst_gold_rgba;
cv::Mat dst_gold_gray;
using FilterTest::SetUpTestCase;
using FilterTest::TearDownTestCase;
virtual void SetUp()
{
Mat kernel(Mat::ones(3, 3, CV_8U));
devInfo = GetParam();
cv::Mat cpuRes;
cv::dilate(img, cpuRes, kernel);
cv::gpu::setDevice(devInfo.deviceID());
GpuMat gpuRes;
cv::gpu::dilate(GpuMat(img), gpuRes, kernel);
return CheckNorm(cpuRes, gpuRes, Size(3, 3));
kernel = cv::Mat::ones(3, 3, CV_8U);
cv::dilate(img_rgba, dst_gold_rgba, kernel);
cv::dilate(img_gray, dst_gold_gray, kernel);
}
};
////////////////////////////////////////////////////////////////////////////////
// MorphologyEx
class CV_GpuMorphExTest : public CV_GpuNppFilterTest
TEST_P(Dilate, Accuracy)
{
public:
CV_GpuMorphExTest() : CV_GpuNppFilterTest( "GPU-NppMorphologyEx", "morphologyEx" ) {}
ASSERT_TRUE(!img_rgba.empty() && !img_gray.empty());
protected:
virtual int test(const Mat& img)
PRINT_PARAM(devInfo);
cv::Mat dst_rgba;
cv::Mat dst_gray;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_dst_rgba;
cv::gpu::GpuMat dev_dst_gray;
cv::gpu::dilate(cv::gpu::GpuMat(img_rgba), dev_dst_rgba, kernel);
cv::gpu::dilate(cv::gpu::GpuMat(img_gray), dev_dst_gray, kernel);
dev_dst_rgba.download(dst_rgba);
dev_dst_gray.download(dst_gray);
);
EXPECT_MAT_NEAR_KSIZE(dst_gold_rgba, dst_rgba, 3, 0.0);
EXPECT_MAT_NEAR_KSIZE(dst_gold_gray, dst_gray, 3, 0.0);
}
INSTANTIATE_TEST_CASE_P(Filter, Dilate, testing::ValuesIn(devices()));
/////////////////////////////////////////////////////////////////////////////////////////////////
// morphEx
static const int morphOps[] = {cv::MORPH_OPEN, CV_MOP_CLOSE, CV_MOP_GRADIENT, CV_MOP_TOPHAT, CV_MOP_BLACKHAT};
static const char* morphOps_str[] = {"MORPH_OPEN", "MOP_CLOSE", "MOP_GRADIENT", "MOP_TOPHAT", "MOP_BLACKHAT"};
struct MorphEx : FilterTest, testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int> >
{
cv::gpu::DeviceInfo devInfo;
int morphOpsIdx;
cv::Mat kernel;
cv::Mat dst_gold_rgba;
cv::Mat dst_gold_gray;
using FilterTest::SetUpTestCase;
using FilterTest::TearDownTestCase;
virtual void SetUp()
{
static int ops[] = { MORPH_OPEN, CV_MOP_CLOSE, CV_MOP_GRADIENT, CV_MOP_TOPHAT, CV_MOP_BLACKHAT};
const char *names[] = { "MORPH_OPEN", "CV_MOP_CLOSE", "CV_MOP_GRADIENT", "CV_MOP_TOPHAT", "CV_MOP_BLACKHAT"};
int num = sizeof(ops)/sizeof(ops[0]);
devInfo = std::tr1::get<0>(GetParam());
morphOpsIdx = std::tr1::get<1>(GetParam());
GpuMat kernel(Mat::ones(3, 3, CV_8U));
cv::gpu::setDevice(devInfo.deviceID());
int res = cvtest::TS::OK;
kernel = cv::Mat::ones(3, 3, CV_8U);
for(int i = 0; i < num; ++i)
{
ts->printf(cvtest::TS::LOG, "Tesing %s\n", names[i]);
cv::Mat cpuRes;
cv::morphologyEx(img, cpuRes, ops[i], (Mat)kernel);
GpuMat gpuRes;
cv::gpu::morphologyEx(GpuMat(img), gpuRes, ops[i], kernel);
if (cvtest::TS::OK != CheckNorm(cpuRes, gpuRes, Size(4, 4)))
res = cvtest::TS::FAIL_GENERIC;
}
return res;
cv::morphologyEx(img_rgba, dst_gold_rgba, morphOps[morphOpsIdx], kernel);
cv::morphologyEx(img_gray, dst_gold_gray, morphOps[morphOpsIdx], kernel);
}
};
TEST_P(MorphEx, Accuracy)
{
ASSERT_TRUE(!img_rgba.empty() && !img_gray.empty());
TEST(blur, accuracy) { CV_GpuNppImageBlurTest test; test.safe_run(); }
TEST(sobel, accuracy) { CV_GpuNppImageSobelTest test; test.safe_run(); }
TEST(scharr, accuracy) { CV_GpuNppImageScharrTest test; test.safe_run(); }
TEST(gaussianBlur, accuracy) { CV_GpuNppImageGaussianBlurTest test; test.safe_run(); }
TEST(laplcaian, accuracy) { CV_GpuNppImageLaplacianTest test; test.safe_run(); }
TEST(erode, accuracy) { CV_GpuErodeTest test; test.safe_run(); }
TEST(dilate, accuracy) { CV_GpuDilateTest test; test.safe_run(); }
TEST(morphEx, accuracy) { CV_GpuMorphExTest test; test.safe_run(); }
const char* morphOpStr = morphOps_str[morphOpsIdx];
PRINT_PARAM(devInfo);
PRINT_PARAM(morphOpStr);
cv::Mat dst_rgba;
cv::Mat dst_gray;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_dst_rgba;
cv::gpu::GpuMat dev_dst_gray;
cv::gpu::morphologyEx(cv::gpu::GpuMat(img_rgba), dev_dst_rgba, morphOps[morphOpsIdx], cv::gpu::GpuMat(kernel));
cv::gpu::morphologyEx(cv::gpu::GpuMat(img_gray), dev_dst_gray, morphOps[morphOpsIdx], cv::gpu::GpuMat(kernel));
dev_dst_rgba.download(dst_rgba);
dev_dst_gray.download(dst_gray);
);
EXPECT_MAT_NEAR_KSIZE(dst_gold_rgba, dst_rgba, 4, 0.0);
EXPECT_MAT_NEAR_KSIZE(dst_gold_gray, dst_gray, 4, 0.0);
}
INSTANTIATE_TEST_CASE_P(Filter, MorphEx, testing::Combine(
testing::ValuesIn(devices()),
testing::Range(0, 5)));
#endif // HAVE_CUDA

154
modules/gpu/test/test_gpu_base.cpp Normal file
View File

@ -0,0 +1,154 @@
/*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.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
bool supportFeature(const cv::gpu::DeviceInfo& info, cv::gpu::FeatureSet feature)
{
return cv::gpu::TargetArchs::builtWith(feature) && info.supports(feature);
}
const std::vector<cv::gpu::DeviceInfo>& devices()
{
static std::vector<cv::gpu::DeviceInfo> devs;
static bool first = true;
if (first)
{
int deviceCount = cv::gpu::getCudaEnabledDeviceCount();
devs.reserve(deviceCount);
for (int i = 0; i < deviceCount; ++i)
{
cv::gpu::DeviceInfo info(i);
if (info.isCompatible())
devs.push_back(info);
}
first = false;
}
return devs;
}
std::vector<cv::gpu::DeviceInfo> devices(cv::gpu::FeatureSet feature)
{
const std::vector<cv::gpu::DeviceInfo>& d = devices();
std::vector<cv::gpu::DeviceInfo> devs_filtered;
if (cv::gpu::TargetArchs::builtWith(feature))
{
devs_filtered.reserve(d.size());
for (size_t i = 0, size = d.size(); i < size; ++i)
{
const cv::gpu::DeviceInfo& info = d[i];
if (info.supports(feature))
devs_filtered.push_back(info);
}
}
return devs_filtered;
}
std::vector<int> types(int depth_start, int depth_end, int cn_start, int cn_end)
{
std::vector<int> v;
v.reserve((depth_end - depth_start + 1) * (cn_end - cn_start + 1));
for (int depth = depth_start; depth <= depth_end; ++depth)
{
for (int cn = cn_start; cn <= cn_end; ++cn)
{
v.push_back(CV_MAKETYPE(depth, cn));
}
}
return v;
}
const std::vector<int>& all_types()
{
static std::vector<int> v = types(CV_8U, CV_64F, 1, 4);
return v;
}
cv::Mat readImage(const std::string& fileName, int flags)
{
return cv::imread(std::string(cvtest::TS::ptr()->get_data_path()) + fileName, flags);
}
double checkNorm(const cv::Mat& m1, const cv::Mat& m2)
{
return cv::norm(m1, m2, cv::NORM_INF);
}
double checkSimilarity(const cv::Mat& m1, const cv::Mat& m2)
{
cv::Mat diff;
cv::matchTemplate(m1, m2, diff, CV_TM_CCORR_NORMED);
return std::abs(diff.at<float>(0, 0) - 1.f);
}
namespace cv
{
std::ostream& operator << (std::ostream& os, const Size& sz)
{
return os << sz.width << "x" << sz.height;
}
std::ostream& operator << (std::ostream& os, const Scalar& s)
{
return os << "[" << s[0] << ", " << s[1] << ", " << s[2] << ", " << s[3] << "]";
}
namespace gpu
{
std::ostream& operator << (std::ostream& os, const DeviceInfo& info)
{
return os << info.name();
}
}
}

103
modules/gpu/test/test_gpu_base.hpp Normal file
View File

@ -0,0 +1,103 @@
/*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.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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*/
#ifndef __OPENCV_TEST_GPU_BASE_HPP__
#define __OPENCV_TEST_GPU_BASE_HPP__
//! return true if device supports specified feature and gpu module was built with support the feature.
bool supportFeature(const cv::gpu::DeviceInfo& info, cv::gpu::FeatureSet feature);
//! return all devices compatible with current gpu module build.
const std::vector<cv::gpu::DeviceInfo>& devices();
//! return all devices compatible with current gpu module build which support specified feature.
std::vector<cv::gpu::DeviceInfo> devices(cv::gpu::FeatureSet feature);
//! return vector with types from specified range.
std::vector<int> types(int depth_start, int depth_end, int cn_start, int cn_end);
//! return vector with all types (depth: CV_8U-CV_64F, channels: 1-4).
const std::vector<int>& all_types();
//! read image from testdata folder.
cv::Mat readImage(const std::string& fileName, int flags = CV_LOAD_IMAGE_COLOR);
double checkNorm(const cv::Mat& m1, const cv::Mat& m2);
double checkSimilarity(const cv::Mat& m1, const cv::Mat& m2);
#define OSTR_NAME(suf) ostr_ ## suf
#define PRINT_PARAM(name) \
std::ostringstream OSTR_NAME(name); \
OSTR_NAME(name) << # name << ": " << name; \
SCOPED_TRACE(OSTR_NAME(name).str());
#define PRINT_TYPE(type) \
std::ostringstream OSTR_NAME(type); \
OSTR_NAME(type) << # type << ": " << cvtest::getTypeName(type) << "c" << CV_MAT_CN(type); \
SCOPED_TRACE(OSTR_NAME(type).str());
#define EXPECT_MAT_NEAR(mat1, mat2, eps) \
{ \
ASSERT_EQ(mat1.type(), mat2.type()); \
ASSERT_EQ(mat1.size(), mat2.size()); \
EXPECT_LE(checkNorm(mat1, mat2), eps); \
}
#define EXPECT_MAT_SIMILAR(mat1, mat2, eps) \
{ \
ASSERT_EQ(mat1.type(), mat2.type()); \
ASSERT_EQ(mat1.size(), mat2.size()); \
EXPECT_LE(checkSimilarity(mat1, mat2), eps); \
}
//! for gtest ASSERT
namespace cv
{
std::ostream& operator << (std::ostream& os, const Size& sz);
std::ostream& operator << (std::ostream& os, const Scalar& s);
namespace gpu
{
std::ostream& operator << (std::ostream& os, const DeviceInfo& info);
}
}
#endif // __OPENCV_TEST_GPU_BASE_HPP__

143
modules/gpu/test/test_hog.cpp
View File

@ -40,32 +40,24 @@
//M*/
#include "test_precomp.hpp"
#include <fstream>
using namespace std;
#ifdef HAVE_CUDA
//#define DUMP
#define CHECK(pred, err) if (!(pred)) { \
ts->printf(cvtest::TS::CONSOLE, "Fail: \"%s\" at line: %d\n", #pred, __LINE__); \
ts->set_failed_test_info(err); \
return; }
struct CV_GpuHogDetectTestRunner: cv::gpu::HOGDescriptor
struct CV_GpuHogDetectTestRunner : cv::gpu::HOGDescriptor
{
CV_GpuHogDetectTestRunner(cvtest::TS* ts_): ts(ts_) {}
void run(int)
void run()
{
cv::Mat img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/road.png");
CHECK(!img_rgb.empty(), cvtest::TS::FAIL_MISSING_TEST_DATA);
cv::Mat img_rgb = readImage("hog/road.png");
ASSERT_TRUE(!img_rgb.empty());
#ifdef DUMP
f.open((std::string(ts->get_data_path()) + "hog/expected_output.bin").c_str(), std::ios_base::binary);
CHECK(f.is_open(), cvtest::TS::FAIL_GENERIC);
f.open((std::string(cvtest::TS::ptr()->get_data_path()) + "hog/expected_output.bin").c_str(), std::ios_base::binary);
ASSERT_TRUE(f.is_open());
#else
f.open((std::string(ts->get_data_path()) + "hog/expected_output.bin").c_str(), std::ios_base::binary);
CHECK(f.is_open(), cvtest::TS::FAIL_MISSING_TEST_DATA);
f.open((std::string(cvtest::TS::ptr()->get_data_path()) + "hog/expected_output.bin").c_str(), std::ios_base::binary);
ASSERT_TRUE(f.is_open());
#endif
// Test on color image
@ -78,7 +70,6 @@ struct CV_GpuHogDetectTestRunner: cv::gpu::HOGDescriptor
test(img);
f.close();
}
#ifdef DUMP
@ -107,24 +98,24 @@ struct CV_GpuHogDetectTestRunner: cv::gpu::HOGDescriptor
f.read((char*)&rows, sizeof(rows));
f.read((char*)&cols, sizeof(cols));
CHECK(rows == block_hists.rows, cvtest::TS::FAIL_INVALID_OUTPUT);
CHECK(cols == block_hists.cols, cvtest::TS::FAIL_INVALID_OUTPUT);
ASSERT_EQ(rows, block_hists.rows);
ASSERT_EQ(cols, block_hists.cols);
for (int i = 0; i < block_hists.rows; ++i)
{
for (int j = 0; j < block_hists.cols; ++j)
{
float val;
f.read((char*)&val, sizeof(val));
CHECK(fabs(val - block_hists.at<float>(i, j)) < 1e-3f, cvtest::TS::FAIL_INVALID_OUTPUT);
ASSERT_NEAR(val, block_hists.at<float>(i, j), 1e-3);
}
}
f.read((char*)&nlocations, sizeof(nlocations));
CHECK(nlocations == static_cast<int>(locations.size()), cvtest::TS::FAIL_INVALID_OUTPUT);
ASSERT_EQ(nlocations, static_cast<int>(locations.size()));
for (int i = 0; i < nlocations; ++i)
{
cv::Point location;
f.read((char*)&location, sizeof(location));
CHECK(location == locations[i], cvtest::TS::FAIL_INVALID_OUTPUT);
ASSERT_EQ(location, locations[i]);
}
}
#endif
@ -176,39 +167,47 @@ struct CV_GpuHogDetectTestRunner: cv::gpu::HOGDescriptor
#else
std::ifstream f;
#endif
cvtest::TS* ts;
};
struct CV_GpuHogDetectTest: cvtest::BaseTest
struct HogDetect : testing::TestWithParam<cv::gpu::DeviceInfo>
{
CV_GpuHogDetectTest() {}
void run(int i)
cv::gpu::DeviceInfo devInfo;
virtual void SetUp()
{
CV_GpuHogDetectTestRunner runner(ts);
runner.run(i);
devInfo = GetParam();
cv::gpu::setDevice(devInfo.deviceID());
}
};
TEST(HOG, detect_accuracy) { CV_GpuHogDetectTest test; test.safe_run(); }
struct CV_GpuHogGetDescriptorsTestRunner: cv::gpu::HOGDescriptor
TEST_P(HogDetect, Accuracy)
{
CV_GpuHogGetDescriptorsTestRunner(cvtest::TS* ts_): HOGDescriptor(cv::Size(64, 128)), ts(ts_) {}
PRINT_PARAM(devInfo);
void run(int)
ASSERT_NO_THROW(
CV_GpuHogDetectTestRunner runner;
runner.run();
);
}
INSTANTIATE_TEST_CASE_P(HOG, HogDetect, testing::ValuesIn(devices()));
struct CV_GpuHogGetDescriptorsTestRunner : cv::gpu::HOGDescriptor
{
CV_GpuHogGetDescriptorsTestRunner(): cv::gpu::HOGDescriptor(cv::Size(64, 128)) {}
void run()
{
// Load image (e.g. train data, composed from windows)
cv::Mat img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/train_data.png");
CHECK(!img_rgb.empty(), cvtest::TS::FAIL_MISSING_TEST_DATA);
cv::Mat img_rgb = readImage("hog/train_data.png");
ASSERT_TRUE(!img_rgb.empty());
// Convert to C4
cv::Mat img;
cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
cv::gpu::GpuMat d_img(img);
cv::gpu::GpuMat d_img(img);
// Convert train images into feature vectors (train table)
cv::gpu::GpuMat descriptors, descriptors_by_cols;
@ -223,7 +222,7 @@ struct CV_GpuHogGetDescriptorsTestRunner: cv::gpu::HOGDescriptor
block_hist_size = 36;
cv::Size descr_size_expected = cv::Size(blocks_per_win_x * blocks_per_win_y * block_hist_size,
wins_per_img_x * wins_per_img_y);
CHECK(descriptors.size() == descr_size_expected, cvtest::TS::FAIL_INVALID_OUTPUT);
ASSERT_EQ(descr_size_expected, descriptors.size());
// Check both formats of output descriptors are handled correctly
cv::Mat dr(descriptors);
@ -235,8 +234,8 @@ struct CV_GpuHogGetDescriptorsTestRunner: cv::gpu::HOGDescriptor
for (int y = 0; y < blocks_per_win_y; ++y)
for (int x = 0; x < blocks_per_win_x; ++x)
for (int k = 0; k < block_hist_size; ++k)
CHECK(l[(y * blocks_per_win_x + x) * block_hist_size + k] ==
r[(x * blocks_per_win_y + y) * block_hist_size + k], cvtest::TS::FAIL_INVALID_OUTPUT);
ASSERT_EQ(l[(y * blocks_per_win_x + x) * block_hist_size + k],
r[(x * blocks_per_win_y + y) * block_hist_size + k]);
}
/* Now we want to extract the same feature vectors, but from single images. NOTE: results will
@ -244,39 +243,39 @@ struct CV_GpuHogGetDescriptorsTestRunner: cv::gpu::HOGDescriptor
wont't call getDescriptors and will use computeBlockHistograms instead of. computeBlockHistograms
works good, it can be checked in the gpu_hog sample */
img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/positive1.png");
CHECK(!img_rgb.empty(), cvtest::TS::FAIL_MISSING_TEST_DATA);
img_rgb = readImage("hog/positive1.png");
ASSERT_TRUE(!img_rgb.empty());
cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
computeBlockHistograms(cv::gpu::GpuMat(img));
// Everything is fine with interpolation for left top subimage
CHECK(cv::norm((cv::Mat)block_hists, (cv::Mat)descriptors.rowRange(0, 1)) == 0.f, cvtest::TS::FAIL_INVALID_OUTPUT);
ASSERT_EQ(0.0, cv::norm((cv::Mat)block_hists, (cv::Mat)descriptors.rowRange(0, 1)));
img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/positive2.png");
CHECK(!img_rgb.empty(), cvtest::TS::FAIL_MISSING_TEST_DATA);
img_rgb = readImage("hog/positive2.png");
ASSERT_TRUE(!img_rgb.empty());
cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
computeBlockHistograms(cv::gpu::GpuMat(img));
compare_inner_parts(block_hists, descriptors.rowRange(1, 2));
img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/negative1.png");
CHECK(!img_rgb.empty(), cvtest::TS::FAIL_MISSING_TEST_DATA);
img_rgb = readImage("hog/negative1.png");
ASSERT_TRUE(!img_rgb.empty());
cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
computeBlockHistograms(cv::gpu::GpuMat(img));
compare_inner_parts(block_hists, descriptors.rowRange(2, 3));
img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/negative2.png");
CHECK(!img_rgb.empty(), cvtest::TS::FAIL_MISSING_TEST_DATA);
img_rgb = readImage("hog/negative2.png");
ASSERT_TRUE(!img_rgb.empty());
cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
computeBlockHistograms(cv::gpu::GpuMat(img));
compare_inner_parts(block_hists, descriptors.rowRange(3, 4));
img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/positive3.png");
CHECK(!img_rgb.empty(), cvtest::TS::FAIL_MISSING_TEST_DATA);
img_rgb = readImage("hog/positive3.png");
ASSERT_TRUE(!img_rgb.empty());
cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
computeBlockHistograms(cv::gpu::GpuMat(img));
compare_inner_parts(block_hists, descriptors.rowRange(4, 5));
img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/negative3.png");
CHECK(!img_rgb.empty(), cvtest::TS::FAIL_MISSING_TEST_DATA);
img_rgb = readImage("hog/negative3.png");
ASSERT_TRUE(!img_rgb.empty());
cv::cvtColor(img_rgb, img, CV_BGR2BGRA);
computeBlockHistograms(cv::gpu::GpuMat(img));
compare_inner_parts(block_hists, descriptors.rowRange(5, 6));
@ -291,7 +290,7 @@ struct CV_GpuHogGetDescriptorsTestRunner: cv::gpu::HOGDescriptor
{
float a = d1.at<float>(0, (i * blocks_per_win_x + j) * block_hist_size);
float b = d2.at<float>(0, (i * blocks_per_win_x + j) * block_hist_size);
CHECK(a == b, cvtest::TS::FAIL_INVALID_OUTPUT)
ASSERT_FLOAT_EQ(a, b);
}
}
@ -300,20 +299,30 @@ struct CV_GpuHogGetDescriptorsTestRunner: cv::gpu::HOGDescriptor
int blocks_per_win_x;
int blocks_per_win_y;
int block_hist_size;
cvtest::TS* ts;
};
struct CV_GpuHogGetDescriptorsTest: cvtest::BaseTest
struct HogGetDescriptors : testing::TestWithParam<cv::gpu::DeviceInfo>
{
CV_GpuHogGetDescriptorsTest() {}
void run(int i)
cv::gpu::DeviceInfo devInfo;
virtual void SetUp()
{
CV_GpuHogGetDescriptorsTestRunner runner(ts);
runner.run(i);
devInfo = GetParam();
cv::gpu::setDevice(devInfo.deviceID());
}
};
TEST(HOG, descriptors_accuracy) { CV_GpuHogGetDescriptorsTest test; test.safe_run(); }
TEST_P(HogGetDescriptors, Accuracy)
{
PRINT_PARAM(devInfo);
ASSERT_NO_THROW(
CV_GpuHogGetDescriptorsTestRunner runner;
runner.run();
);
}
INSTANTIATE_TEST_CASE_P(HOG, HogGetDescriptors, testing::ValuesIn(devices()));
#endif // HAVE_CUDA

2137
modules/gpu/test/test_imgproc.cpp Normal file
View File

File diff suppressed because it is too large Load Diff

966
modules/gpu/test/test_imgproc_gpu.cpp
View File

@ -1,966 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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 <cmath>
#include <limits>
#include "test_precomp.hpp"
using namespace cv;
using namespace std;
using namespace gpu;
class CV_GpuImageProcTest : public cvtest::BaseTest
{
public:
virtual ~CV_GpuImageProcTest() {}
protected:
void run(int);
int test8UC1 (const Mat& img);
int test8UC4 (const Mat& img);
int test32SC1(const Mat& img);
int test32FC1(const Mat& img);
virtual int test(const Mat& img) = 0;
int CheckNorm(const Mat& m1, const Mat& m2);
// Checks whether two images are similar enough using normalized
// cross-correlation as an error measure
int CheckSimilarity(const Mat& m1, const Mat& m2, float max_err=1e-3f);
};
int CV_GpuImageProcTest::test8UC1(const Mat& img)
{
cv::Mat img_C1;
cvtColor(img, img_C1, CV_BGR2GRAY);
return test(img_C1);
}
int CV_GpuImageProcTest::test8UC4(const Mat& img)
{
cv::Mat img_C4;
cvtColor(img, img_C4, CV_BGR2BGRA);
return test(img_C4);
}
int CV_GpuImageProcTest::test32SC1(const Mat& img)
{
cv::Mat img_C1;
cvtColor(img, img_C1, CV_BGR2GRAY);
img_C1.convertTo(img_C1, CV_32S);
return test(img_C1);
}
int CV_GpuImageProcTest::test32FC1(const Mat& img)
{
cv::Mat temp, img_C1;
img.convertTo(temp, CV_32F, 1.f / 255.f);
cvtColor(temp, img_C1, CV_BGR2GRAY);
return test(img_C1);
}
int CV_GpuImageProcTest::CheckNorm(const Mat& m1, const Mat& m2)
{
double ret = norm(m1, m2, NORM_INF);
if (ret < std::numeric_limits<double>::epsilon())
{
return cvtest::TS::OK;
}
else
{
ts->printf(cvtest::TS::LOG, "Norm: %f\n", ret);
return cvtest::TS::FAIL_GENERIC;
}
}
int CV_GpuImageProcTest::CheckSimilarity(const Mat& m1, const Mat& m2, float max_err)
{
Mat diff;
cv::matchTemplate(m1, m2, diff, CV_TM_CCORR_NORMED);
float err = abs(diff.at<float>(0, 0) - 1.f);
if (err > max_err)
return cvtest::TS::FAIL_INVALID_OUTPUT;
return cvtest::TS::OK;
}
void CV_GpuImageProcTest::run( int )
{
//load image
cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");
if (img.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
int testResult = cvtest::TS::OK;
//run tests
ts->printf(cvtest::TS::LOG, "\n========Start test 8UC1========\n");
if (test8UC1(img) == cvtest::TS::OK)
ts->printf(cvtest::TS::LOG, "SUCCESS\n");
else
{
ts->printf(cvtest::TS::LOG, "FAIL\n");
testResult = cvtest::TS::FAIL_GENERIC;
}
ts->printf(cvtest::TS::LOG, "\n========Start test 8UC4========\n");
if (test8UC4(img) == cvtest::TS::OK)
ts->printf(cvtest::TS::LOG, "SUCCESS\n");
else
{
ts->printf(cvtest::TS::LOG, "FAIL\n");
testResult = cvtest::TS::FAIL_GENERIC;
}
ts->printf(cvtest::TS::LOG, "\n========Start test 32SC1========\n");
if (test32SC1(img) == cvtest::TS::OK)
ts->printf(cvtest::TS::LOG, "SUCCESS\n");
else
{
ts->printf(cvtest::TS::LOG, "FAIL\n");
testResult = cvtest::TS::FAIL_GENERIC;
}
ts->printf(cvtest::TS::LOG, "\n========Start test 32FC1========\n");
if (test32FC1(img) == cvtest::TS::OK)
ts->printf(cvtest::TS::LOG, "SUCCESS\n");
else
{
ts->printf(cvtest::TS::LOG, "FAIL\n");
testResult = cvtest::TS::FAIL_GENERIC;
}
ts->set_failed_test_info(testResult);
}
////////////////////////////////////////////////////////////////////////////////
// threshold
struct CV_GpuImageThresholdTest : public CV_GpuImageProcTest
{
public:
CV_GpuImageThresholdTest() {}
int test(const Mat& img)
{
if (img.type() != CV_8UC1 && img.type() != CV_32FC1)
{
ts->printf(cvtest::TS::LOG, "\nUnsupported type\n");
return cvtest::TS::OK;
}
const double maxVal = img.type() == CV_8UC1 ? 255 : 1.0;
cv::RNG& rng = ts->get_rng();
int res = cvtest::TS::OK;
for (int type = THRESH_BINARY; type <= THRESH_TOZERO_INV; ++type)
{
const double thresh = rng.uniform(0.0, maxVal);
cv::Mat cpuRes;
cv::threshold(img, cpuRes, thresh, maxVal, type);
GpuMat gpu1(img);
GpuMat gpuRes;
cv::gpu::threshold(gpu1, gpuRes, thresh, maxVal, type);
if (CheckNorm(cpuRes, gpuRes) != cvtest::TS::OK)
res = cvtest::TS::FAIL_GENERIC;
}
return res;
}
};
////////////////////////////////////////////////////////////////////////////////
// resize
struct CV_GpuNppImageResizeTest : public CV_GpuImageProcTest
{
CV_GpuNppImageResizeTest() {}
int test(const Mat& img)
{
if (img.type() != CV_8UC1 && img.type() != CV_8UC4)
{
ts->printf(cvtest::TS::LOG, "Unsupported type\n");
return cvtest::TS::OK;
}
int interpolations[] = {INTER_NEAREST, INTER_LINEAR, /*INTER_CUBIC,*/ /*INTER_LANCZOS4*/};
const char* interpolations_str[] = {"INTER_NEAREST", "INTER_LINEAR", /*"INTER_CUBIC",*/ /*"INTER_LANCZOS4"*/};
int interpolations_num = sizeof(interpolations) / sizeof(int);
int test_res = cvtest::TS::OK;
for (int i = 0; i < interpolations_num; ++i)
{
ts->printf(cvtest::TS::LOG, "Interpolation: %s\n", interpolations_str[i]);
Mat cpu_res1, cpu_res2;
cv::resize(img, cpu_res1, Size(), 2.0, 2.0, interpolations[i]);
cv::resize(cpu_res1, cpu_res2, Size(), 0.5, 0.5, interpolations[i]);
GpuMat gpu1(img), gpu_res1, gpu_res2;
cv::gpu::resize(gpu1, gpu_res1, Size(), 2.0, 2.0, interpolations[i]);
cv::gpu::resize(gpu_res1, gpu_res2, Size(), 0.5, 0.5, interpolations[i]);
if (CheckSimilarity(cpu_res2, gpu_res2) != cvtest::TS::OK)
test_res = cvtest::TS::FAIL_GENERIC;
}
return test_res;
}
};
////////////////////////////////////////////////////////////////////////////////
// copyMakeBorder
struct CV_GpuNppImageCopyMakeBorderTest : public CV_GpuImageProcTest
{
CV_GpuNppImageCopyMakeBorderTest() {}
int test(const Mat& img)
{
if (img.type() != CV_8UC1 && img.type() != CV_8UC4 && img.type() != CV_32SC1)
{
ts->printf(cvtest::TS::LOG, "\nUnsupported type\n");
return cvtest::TS::OK;
}
cv::RNG& rng = ts->get_rng();
int top = rng.uniform(1, 10);
int botton = rng.uniform(1, 10);
int left = rng.uniform(1, 10);
int right = rng.uniform(1, 10);
cv::Scalar val(rng.uniform(0, 255), rng.uniform(0, 255), rng.uniform(0, 255), rng.uniform(0, 255));
Mat cpudst;
cv::copyMakeBorder(img, cpudst, top, botton, left, right, BORDER_CONSTANT, val);
GpuMat gpu1(img);
GpuMat gpudst;
cv::gpu::copyMakeBorder(gpu1, gpudst, top, botton, left, right, val);
return CheckNorm(cpudst, gpudst);
}
};
////////////////////////////////////////////////////////////////////////////////
// warpAffine
struct CV_GpuNppImageWarpAffineTest : public CV_GpuImageProcTest
{
CV_GpuNppImageWarpAffineTest() {}
int test(const Mat& img)
{
if (img.type() == CV_32SC1)
{
ts->printf(cvtest::TS::LOG, "\nUnsupported type\n");
return cvtest::TS::OK;
}
static double reflect[2][3] = { {-1, 0, 0},
{ 0, -1, 0} };
reflect[0][2] = img.cols;
reflect[1][2] = img.rows;
Mat M(2, 3, CV_64F, (void*)reflect);
int flags[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_NEAREST | WARP_INVERSE_MAP, INTER_LINEAR | WARP_INVERSE_MAP, INTER_CUBIC | WARP_INVERSE_MAP};
const char* flags_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_NEAREST | WARP_INVERSE_MAP", "INTER_LINEAR | WARP_INVERSE_MAP", "INTER_CUBIC | WARP_INVERSE_MAP"};
int flags_num = sizeof(flags) / sizeof(int);
int test_res = cvtest::TS::OK;
for (int i = 0; i < flags_num; ++i)
{
ts->printf(cvtest::TS::LOG, "\nFlags: %s\n", flags_str[i]);
Mat cpudst;
cv::warpAffine(img, cpudst, M, img.size(), flags[i]);
GpuMat gpu1(img);
GpuMat gpudst;
cv::gpu::warpAffine(gpu1, gpudst, M, gpu1.size(), flags[i]);
// Check inner parts (ignoring 1 pixel width border)
if (CheckSimilarity(cpudst.rowRange(1, cpudst.rows - 1).colRange(1, cpudst.cols - 1),
gpudst.rowRange(1, gpudst.rows - 1).colRange(1, gpudst.cols - 1)) != cvtest::TS::OK)
test_res = cvtest::TS::FAIL_GENERIC;
}
return test_res;
}
};
////////////////////////////////////////////////////////////////////////////////
// warpPerspective
struct CV_GpuNppImageWarpPerspectiveTest : public CV_GpuImageProcTest
{
CV_GpuNppImageWarpPerspectiveTest() {}
int test(const Mat& img)
{
if (img.type() == CV_32SC1)
{
ts->printf(cvtest::TS::LOG, "\nUnsupported type\n");
return cvtest::TS::OK;
}
static double reflect[3][3] = { { -1, 0, 0},
{ 0, -1, 0},
{ 0, 0, 1 }};
reflect[0][2] = img.cols;
reflect[1][2] = img.rows;
Mat M(3, 3, CV_64F, (void*)reflect);
int flags[] = {INTER_NEAREST, INTER_LINEAR, INTER_CUBIC, INTER_NEAREST | WARP_INVERSE_MAP, INTER_LINEAR | WARP_INVERSE_MAP, INTER_CUBIC | WARP_INVERSE_MAP};
const char* flags_str[] = {"INTER_NEAREST", "INTER_LINEAR", "INTER_CUBIC", "INTER_NEAREST | WARP_INVERSE_MAP", "INTER_LINEAR | WARP_INVERSE_MAP", "INTER_CUBIC | WARP_INVERSE_MAP"};
int flags_num = sizeof(flags) / sizeof(int);
int test_res = cvtest::TS::OK;
for (int i = 0; i < flags_num; ++i)
{
ts->printf(cvtest::TS::LOG, "\nFlags: %s\n", flags_str[i]);
Mat cpudst;
cv::warpPerspective(img, cpudst, M, img.size(), flags[i]);
GpuMat gpu1(img);
GpuMat gpudst;
cv::gpu::warpPerspective(gpu1, gpudst, M, gpu1.size(), flags[i]);
// Check inner parts (ignoring 1 pixel width border)
if (CheckSimilarity(cpudst.rowRange(1, cpudst.rows - 1).colRange(1, cpudst.cols - 1),
gpudst.rowRange(1, gpudst.rows - 1).colRange(1, gpudst.cols - 1)) != cvtest::TS::OK)
test_res = cvtest::TS::FAIL_GENERIC;
}
return test_res;
}
};
////////////////////////////////////////////////////////////////////////////////
// integral
struct CV_GpuNppImageIntegralTest : public CV_GpuImageProcTest
{
CV_GpuNppImageIntegralTest() {}
int test(const Mat& img)
{
if (img.type() != CV_8UC1)
{
ts->printf(cvtest::TS::LOG, "\nUnsupported type\n");
return cvtest::TS::OK;
}
Mat cpusum;
cv::integral(img, cpusum, CV_32S);
GpuMat gpu1(img);
GpuMat gpusum;
cv::gpu::integral(gpu1, gpusum);
return CheckNorm(cpusum, gpusum) == cvtest::TS::OK ? cvtest::TS::OK : cvtest::TS::FAIL_GENERIC;
}
};
////////////////////////////////////////////////////////////////////////////////
// Canny
//struct CV_GpuNppImageCannyTest : public CV_GpuImageProcTest
//{
// CV_GpuNppImageCannyTest() : CV_GpuImageProcTest( "GPU-NppImageCanny", "Canny" ) {}
//
// int test(const Mat& img)
// {
// if (img.type() != CV_8UC1)
// {
// ts->printf(cvtest::TS::LOG, "\nUnsupported type\n");
// return cvtest::TS::OK;
// }
//
// const double threshold1 = 1.0, threshold2 = 10.0;
//
// Mat cpudst;
// cv::Canny(img, cpudst, threshold1, threshold2);
//
// GpuMat gpu1(img);
// GpuMat gpudst;
// cv::gpu::Canny(gpu1, gpudst, threshold1, threshold2);
//
// return CheckNorm(cpudst, gpudst);
// }
//};
////////////////////////////////////////////////////////////////////////////////
// cvtColor
class CV_GpuCvtColorTest : public cvtest::BaseTest
{
public:
CV_GpuCvtColorTest() {}
~CV_GpuCvtColorTest() {};
protected:
void run(int);
int CheckNorm(const Mat& m1, const Mat& m2);
};
int CV_GpuCvtColorTest::CheckNorm(const Mat& m1, const Mat& m2)
{
float max_err = 1e-2f;
Mat diff;
cv::matchTemplate(m1, m2, diff, CV_TM_CCORR_NORMED);
float err = abs(diff.at<float>(0, 0) - 1.f);
if (err > max_err)
return cvtest::TS::FAIL_INVALID_OUTPUT;
return cvtest::TS::OK;
}
void CV_GpuCvtColorTest::run( int )
{
cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");
if (img.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
int testResult = cvtest::TS::OK;
cv::Mat cpuRes;
cv::gpu::GpuMat gpuImg(img), gpuRes;
int codes[] = { CV_BGR2RGB, CV_RGB2BGRA, CV_BGRA2RGB,
CV_RGB2BGR555, CV_BGR5552BGR, CV_BGR2BGR565, CV_BGR5652RGB,
CV_RGB2YCrCb, CV_YCrCb2BGR, CV_BGR2YUV, CV_YUV2RGB,
CV_RGB2XYZ, CV_XYZ2BGR, CV_BGR2XYZ, CV_XYZ2RGB,
CV_RGB2HSV, CV_HSV2BGR, CV_BGR2HSV_FULL, CV_HSV2RGB_FULL,
CV_RGB2HLS, CV_HLS2BGR, CV_BGR2HLS_FULL, CV_HLS2RGB_FULL,
CV_RGB2GRAY, CV_GRAY2BGRA, CV_BGRA2GRAY,
CV_GRAY2BGR555, CV_BGR5552GRAY, CV_GRAY2BGR565, CV_BGR5652GRAY};
const char* codes_str[] = { "CV_BGR2RGB", "CV_RGB2BGRA", "CV_BGRA2RGB",
"CV_RGB2BGR555", "CV_BGR5552BGR", "CV_BGR2BGR565", "CV_BGR5652RGB",
"CV_RGB2YCrCb", "CV_YCrCb2BGR", "CV_BGR2YUV", "CV_YUV2RGB",
"CV_RGB2XYZ", "CV_XYZ2BGR", "CV_BGR2XYZ", "CV_XYZ2RGB",
"CV_RGB2HSV", "CV_HSV2RGB", "CV_BGR2HSV_FULL", "CV_HSV2RGB_FULL",
"CV_RGB2HLS", "CV_HLS2RGB", "CV_BGR2HLS_FULL", "CV_HLS2RGB_FULL",
"CV_RGB2GRAY", "CV_GRAY2BGRA", "CV_BGRA2GRAY",
"CV_GRAY2BGR555", "CV_BGR5552GRAY", "CV_GRAY2BGR565", "CV_BGR5652GRAY"};
int codes_num = sizeof(codes) / sizeof(int);
for (int i = 0; i < codes_num; ++i)
{
ts->printf(cvtest::TS::LOG, "\n%s\n", codes_str[i]);
cv::cvtColor(img, cpuRes, codes[i]);
cv::gpu::cvtColor(gpuImg, gpuRes, codes[i]);
if (CheckNorm(cpuRes, gpuRes) == cvtest::TS::OK)
ts->printf(cvtest::TS::LOG, "\nSUCCESS\n");
else
{
ts->printf(cvtest::TS::LOG, "\nFAIL\n");
testResult = cvtest::TS::FAIL_GENERIC;
}
img = cpuRes;
gpuImg = gpuRes;
}
ts->set_failed_test_info(testResult);
}
////////////////////////////////////////////////////////////////////////////////
// Histograms
class CV_GpuHistogramsTest : public cvtest::BaseTest
{
public:
CV_GpuHistogramsTest() {}
~CV_GpuHistogramsTest() {};
protected:
void run(int);
int CheckNorm(const Mat& m1, const Mat& m2)
{
double ret = norm(m1, m2, NORM_INF);
if (ret < std::numeric_limits<double>::epsilon())
{
return cvtest::TS::OK;
}
else
{
ts->printf(cvtest::TS::LOG, "\nNorm: %f\n", ret);
return cvtest::TS::FAIL_GENERIC;
}
}
};
void CV_GpuHistogramsTest::run( int )
{
//load image
cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");
if (img.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
Mat hsv;
cv::cvtColor(img, hsv, CV_BGR2HSV);
int hbins = 30;
int histSize[] = {hbins};
float hranges[] = {0, 180};
const float* ranges[] = {hranges};
MatND hist;
int channels[] = {0};
calcHist(&hsv, 1, channels, Mat(), hist, 1, histSize, ranges);
GpuMat gpuHsv(hsv);
std::vector<GpuMat> srcs;
cv::gpu::split(gpuHsv, srcs);
GpuMat gpuHist;
histEven(srcs[0], gpuHist, hbins, (int)hranges[0], (int)hranges[1]);
Mat cpuHist = hist;
cpuHist = cpuHist.t();
cpuHist.convertTo(cpuHist, CV_32S);
ts->set_failed_test_info(CheckNorm(cpuHist, gpuHist));
}
////////////////////////////////////////////////////////////////////////
// Corner Harris feature detector
struct CV_GpuCornerHarrisTest: cvtest::BaseTest
{
CV_GpuCornerHarrisTest() {}
void run(int)
{
for (int i = 0; i < 5; ++i)
{
int rows = 25 + rand() % 300, cols = 25 + rand() % 300;
if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, 1 + 2 * (rand() % 4))) return;
if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, -1)) return;
if (!compareToCpuTest(rows, cols, CV_8U, 1 + rand() % 5, 1 + 2 * (rand() % 4))) return;
if (!compareToCpuTest(rows, cols, CV_8U, 1 + rand() % 5, -1)) return;
}
}
bool compareToCpuTest(int rows, int cols, int depth, int blockSize, int apertureSize)
{
RNG rng;
cv::Mat src(rows, cols, depth);
if (depth == CV_32F)
rng.fill(src, RNG::UNIFORM, cv::Scalar(0), cv::Scalar(1));
else if (depth == CV_8U)
rng.fill(src, RNG::UNIFORM, cv::Scalar(0), cv::Scalar(256));
double k = 0.1;
cv::Mat dst_gold;
cv::gpu::GpuMat dst;
cv::Mat dsth;
int borderType;
borderType = BORDER_REFLECT101;
cv::cornerHarris(src, dst_gold, blockSize, apertureSize, k, borderType);
cv::gpu::cornerHarris(cv::gpu::GpuMat(src), dst, blockSize, apertureSize, k, borderType);
dsth = dst;
for (int i = 0; i < dst.rows; ++i)
{
for (int j = 0; j < dst.cols; ++j)
{
float a = dst_gold.at<float>(i, j);
float b = dsth.at<float>(i, j);
if (fabs(a - b) > 1e-3f)
{
ts->printf(cvtest::TS::CONSOLE, "%d %d %f %f %d\n", i, j, a, b, apertureSize);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
};
}
}
borderType = BORDER_REPLICATE;
cv::cornerHarris(src, dst_gold, blockSize, apertureSize, k, borderType);
cv::gpu::cornerHarris(cv::gpu::GpuMat(src), dst, blockSize, apertureSize, k, borderType);
dsth = dst;
for (int i = 0; i < dst.rows; ++i)
{
for (int j = 0; j < dst.cols; ++j)
{
float a = dst_gold.at<float>(i, j);
float b = dsth.at<float>(i, j);
if (fabs(a - b) > 1e-3f)
{
ts->printf(cvtest::TS::CONSOLE, "%d %d %f %f %d\n", i, j, a, b, apertureSize);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
};
}
}
return true;
}
};
////////////////////////////////////////////////////////////////////////
// Corner Min Eigen Val
struct CV_GpuCornerMinEigenValTest: cvtest::BaseTest
{
CV_GpuCornerMinEigenValTest() {}
void run(int)
{
for (int i = 0; i < 3; ++i)
{
int rows = 25 + rand() % 300, cols = 25 + rand() % 300;
if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, -1)) return;
if (!compareToCpuTest(rows, cols, CV_32F, 1 + rand() % 5, 1 + 2 * (rand() % 4))) return;
if (!compareToCpuTest(rows, cols, CV_8U, 1 + rand() % 5, -1)) return;
if (!compareToCpuTest(rows, cols, CV_8U, 1 + rand() % 5, 1 + 2 * (rand() % 4))) return;
}
}
bool compareToCpuTest(int rows, int cols, int depth, int blockSize, int apertureSize)
{
RNG rng;
cv::Mat src(rows, cols, depth);
if (depth == CV_32F)
rng.fill(src, RNG::UNIFORM, cv::Scalar(0), cv::Scalar(1));
else if (depth == CV_8U)
rng.fill(src, RNG::UNIFORM, cv::Scalar(0), cv::Scalar(256));
cv::Mat dst_gold;
cv::gpu::GpuMat dst;
cv::Mat dsth;
int borderType;
borderType = BORDER_REFLECT101;
cv::cornerMinEigenVal(src, dst_gold, blockSize, apertureSize, borderType);
cv::gpu::cornerMinEigenVal(cv::gpu::GpuMat(src), dst, blockSize, apertureSize, borderType);
dsth = dst;
for (int i = 0; i < dst.rows; ++i)
{
for (int j = 0; j < dst.cols; ++j)
{
float a = dst_gold.at<float>(i, j);
float b = dsth.at<float>(i, j);
if (fabs(a - b) > 1e-2f)
{
ts->printf(cvtest::TS::CONSOLE, "%d %d %f %f %d %d\n", i, j, a, b, apertureSize, blockSize);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
};
}
}
borderType = BORDER_REPLICATE;
cv::cornerMinEigenVal(src, dst_gold, blockSize, apertureSize, borderType);
cv::gpu::cornerMinEigenVal(cv::gpu::GpuMat(src), dst, blockSize, apertureSize, borderType);
dsth = dst;
for (int i = 0; i < dst.rows; ++i)
{
for (int j = 0; j < dst.cols; ++j)
{
float a = dst_gold.at<float>(i, j);
float b = dsth.at<float>(i, j);
if (fabs(a - b) > 1e-2f)
{
ts->printf(cvtest::TS::CONSOLE, "%d %d %f %f %d %d\n", i, j, a, b, apertureSize, blockSize);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
};
}
}
return true;
}
};
struct CV_GpuColumnSumTest: cvtest::BaseTest
{
CV_GpuColumnSumTest() {}
void run(int)
{
int cols = 375;
int rows = 1072;
Mat src(rows, cols, CV_32F);
RNG rng(1);
rng.fill(src, RNG::UNIFORM, Scalar(0), Scalar(1));
GpuMat d_dst;
columnSum(GpuMat(src), d_dst);
Mat dst = d_dst;
for (int j = 0; j < src.cols; ++j)
{
float a = src.at<float>(0, j);
float b = dst.at<float>(0, j);
if (fabs(a - b) > 0.5f)
{
ts->printf(cvtest::TS::CONSOLE, "big diff at %d %d: %f %f\n", 0, j, a, b);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
}
for (int i = 1; i < src.rows; ++i)
{
for (int j = 0; j < src.cols; ++j)
{
float a = src.at<float>(i, j) += src.at<float>(i - 1, j);
float b = dst.at<float>(i, j);
if (fabs(a - b) > 0.5f)
{
ts->printf(cvtest::TS::CONSOLE, "big diff at %d %d: %f %f\n", i, j, a, b);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
}
}
}
};
struct CV_GpuNormTest : cvtest::BaseTest
{
CV_GpuNormTest() {}
void run(int)
{
RNG rng(0);
int rows = rng.uniform(1, 500);
int cols = rng.uniform(1, 500);
for (int cn = 1; cn <= 4; ++cn)
{
test(NORM_L1, rows, cols, CV_8U, cn, Scalar::all(0), Scalar::all(10));
test(NORM_L1, rows, cols, CV_8S, cn, Scalar::all(-10), Scalar::all(10));
test(NORM_L1, rows, cols, CV_16U, cn, Scalar::all(0), Scalar::all(10));
test(NORM_L1, rows, cols, CV_16S, cn, Scalar::all(-10), Scalar::all(10));
test(NORM_L1, rows, cols, CV_32S, cn, Scalar::all(-10), Scalar::all(10));
test(NORM_L1, rows, cols, CV_32F, cn, Scalar::all(0), Scalar::all(1));
test(NORM_L2, rows, cols, CV_8U, cn, Scalar::all(0), Scalar::all(10));
test(NORM_L2, rows, cols, CV_8S, cn, Scalar::all(-10), Scalar::all(10));
test(NORM_L2, rows, cols, CV_16U, cn, Scalar::all(0), Scalar::all(10));
test(NORM_L2, rows, cols, CV_16S, cn, Scalar::all(-10), Scalar::all(10));
test(NORM_L2, rows, cols, CV_32S, cn, Scalar::all(-10), Scalar::all(10));
test(NORM_L2, rows, cols, CV_32F, cn, Scalar::all(0), Scalar::all(1));
test(NORM_INF, rows, cols, CV_8U, cn, Scalar::all(0), Scalar::all(10));
test(NORM_INF, rows, cols, CV_8S, cn, Scalar::all(-10), Scalar::all(10));
test(NORM_INF, rows, cols, CV_16U, cn, Scalar::all(0), Scalar::all(10));
test(NORM_INF, rows, cols, CV_16S, cn, Scalar::all(-10), Scalar::all(10));
test(NORM_INF, rows, cols, CV_32S, cn, Scalar::all(-10), Scalar::all(10));
test(NORM_INF, rows, cols, CV_32F, cn, Scalar::all(0), Scalar::all(1));
}
}
void gen(Mat& mat, int rows, int cols, int type, Scalar low, Scalar high)
{
mat.create(rows, cols, type);
RNG rng(0);
rng.fill(mat, RNG::UNIFORM, low, high);
}
void test(int norm_type, int rows, int cols, int depth, int cn, Scalar low, Scalar high)
{
int type = CV_MAKE_TYPE(depth, cn);
Mat src;
gen(src, rows, cols, type, low, high);
double gold = norm(src, norm_type);
double mine = norm(GpuMat(src), norm_type);
if (abs(gold - mine) > 1e-3)
{
ts->printf(cvtest::TS::CONSOLE, "failed test: gold=%f, mine=%f, norm_type=%d, rows=%d, "
"cols=%d, depth=%d, cn=%d\n", gold, mine, norm_type, rows, cols, depth, cn);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
}
}
};
////////////////////////////////////////////////////////////////////////////////
// reprojectImageTo3D
class CV_GpuReprojectImageTo3DTest : public cvtest::BaseTest
{
public:
CV_GpuReprojectImageTo3DTest() {}
protected:
void run(int)
{
Mat disp(320, 240, CV_8UC1);
RNG& rng = ts->get_rng();
rng.fill(disp, RNG::UNIFORM, Scalar(5), Scalar(30));
Mat Q(4, 4, CV_32FC1);
rng.fill(Q, RNG::UNIFORM, Scalar(0.1), Scalar(1));
Mat cpures;
GpuMat gpures;
reprojectImageTo3D(disp, cpures, Q, false);
reprojectImageTo3D(GpuMat(disp), gpures, Q);
Mat temp = gpures;
for (int y = 0; y < cpures.rows; ++y)
{
const Vec3f* cpu_row = cpures.ptr<Vec3f>(y);
const Vec4f* gpu_row = temp.ptr<Vec4f>(y);
for (int x = 0; x < cpures.cols; ++x)
{
Vec3f a = cpu_row[x];
Vec4f b = gpu_row[x];
if (fabs(a[0] - b[0]) > 1e-5 || fabs(a[1] - b[1]) > 1e-5 || fabs(a[2] - b[2]) > 1e-5)
{
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
}
}
}
};
TEST(threshold, accuracy) { CV_GpuImageThresholdTest test; test.safe_run(); }
TEST(resize, accuracy) { CV_GpuNppImageResizeTest test; test.safe_run(); }
TEST(copyMakeBorder, accuracy) { CV_GpuNppImageCopyMakeBorderTest test; test.safe_run(); }
TEST(warpAffine, accuracy) { CV_GpuNppImageWarpAffineTest test; test.safe_run(); }
TEST(warpPerspective, accuracy) { CV_GpuNppImageWarpPerspectiveTest test; test.safe_run(); }
TEST(integral, accuracy) { CV_GpuNppImageIntegralTest test; test.safe_run(); }
TEST(cvtColor, accuracy) { CV_GpuCvtColorTest test; test.safe_run(); }
TEST(histograms, accuracy) { CV_GpuHistogramsTest test; test.safe_run(); }
TEST(cornerHearris, accuracy) { CV_GpuCornerHarrisTest test; test.safe_run(); }
TEST(minEigen, accuracy) { CV_GpuCornerMinEigenValTest test; test.safe_run(); }
TEST(columnSum, accuracy) { CV_GpuColumnSumTest test; test.safe_run(); }
TEST(norm, accuracy) { CV_GpuNormTest test; test.safe_run(); }
TEST(reprojectImageTo3D, accuracy) { CV_GpuReprojectImageTo3DTest test; test.safe_run(); }
TEST(downsample, accuracy_on_8U)
{
RNG& rng = cvtest::TS::ptr()->get_rng();
Size size(200 + cvtest::randInt(rng) % 1000, 200 + cvtest::randInt(rng) % 1000);
Mat src = cvtest::randomMat(rng, size, CV_8U, 0, 255, false);
for (int k = 2; k <= 5; ++k)
{
GpuMat d_dst;
downsample(GpuMat(src), d_dst, k);
Size dst_gold_size((src.cols + k - 1) / k, (src.rows + k - 1) / k);
ASSERT_EQ(dst_gold_size.width, d_dst.cols)
<< "rows=" << size.height << ", cols=" << size.width << ", k=" << k;
ASSERT_EQ(dst_gold_size.height, d_dst.rows)
<< "rows=" << size.height << ", cols=" << size.width << ", k=" << k;
Mat dst = d_dst;
for (int y = 0; y < dst.rows; ++y)
for (int x = 0; x < dst.cols; ++x)
ASSERT_EQ(src.at<uchar>(y * k, x * k), dst.at<uchar>(y, x))
<< "rows=" << size.height << ", cols=" << size.width << ", k=" << k;
}
}
TEST(downsample, accuracy_on_32F)
{
RNG& rng = cvtest::TS::ptr()->get_rng();
Size size(200 + cvtest::randInt(rng) % 1000, 200 + cvtest::randInt(rng) % 1000);
Mat src = cvtest::randomMat(rng, size, CV_32F, 0, 1, false);
for (int k = 2; k <= 5; ++k)
{
GpuMat d_dst;
downsample(GpuMat(src), d_dst, k);
Size dst_gold_size((src.cols + k - 1) / k, (src.rows + k - 1) / k);
ASSERT_EQ(dst_gold_size.width, d_dst.cols)
<< "rows=" << size.height << ", cols=" << size.width << ", k=" << k;
ASSERT_EQ(dst_gold_size.height, d_dst.rows)
<< "rows=" << size.height << ", cols=" << size.width << ", k=" << k;
Mat dst = d_dst;
for (int y = 0; y < dst.rows; ++y)
for (int x = 0; x < dst.cols; ++x)
ASSERT_FLOAT_EQ(src.at<float>(y * k, x * k), dst.at<float>(y, x))
<< "rows=" << size.height << ", cols=" << size.width << ", k=" << k;
}
}

110
modules/gpu/test/test_main.cpp
View File

@ -1,13 +1,107 @@
/*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.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
int main(int argc, char **argv)
#ifdef HAVE_CUDA
void print_info()
{
int deviceCount = cv::gpu::getCudaEnabledDeviceCount();
printf("Found %d CUDA devices\n\n", deviceCount);
for (int i = 0; i < deviceCount; ++i)
{
cv::gpu::DeviceInfo info(i);
printf("Device %d:\n", i);
printf("\tName: %s\n", info.name().c_str());
printf("\tCompute capability version: %d.%d\n", info.majorVersion(), info.minorVersion());
printf("\tTotal memory: %d Mb\n", static_cast<int>(static_cast<int>(info.totalMemory() / 1024.0) / 1024.0));
printf("\tFree memory: %d Mb\n", static_cast<int>(static_cast<int>(info.freeMemory() / 1024.0) / 1024.0));
if (!info.isCompatible())
printf("\tThis device is not compatible with current GPU module build\n");
printf("\n");
}
printf("GPU module was compiled for next GPU archs:\n");
printf("\tBIN:%s\n", CUDA_ARCH_BIN);
printf("\tPTX:%s\n\n", CUDA_ARCH_PTX);
}
enum OutputLevel
{
OutputLevelNone,
OutputLevelCompact,
OutputLevelFull
};
extern OutputLevel nvidiaTestOutputLevel;
int main(int argc, char** argv)
{
cvtest::TS::ptr()->init("gpu");
::testing::InitGoogleTest(&argc, argv);
#ifdef HAVE_CUDA
testing::InitGoogleTest(&argc, argv);
cv::CommandLineParser parser(argc, (const char**)argv);
std::string outputLevel = parser.get<std::string>("nvtest_output_level", "none");
if (outputLevel == "none")
nvidiaTestOutputLevel = OutputLevelNone;
else if (outputLevel == "compact")
nvidiaTestOutputLevel = OutputLevelCompact;
else if (outputLevel == "full")
nvidiaTestOutputLevel = OutputLevelFull;
print_info();
return RUN_ALL_TESTS();
#else
std::cerr << "opencv_test_gpu: OpenCV was compiled without GPU support\n";
return -1;
#endif
}
}
#else // HAVE_CUDA
int main(int argc, char** argv)
{
printf("OpenCV was built without CUDA support\n");
return 0;
}
#endif // HAVE_CUDA

295
modules/gpu/test/test_match_template.cpp
View File

@ -1,295 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other GpuMaterials 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 bpied warranties, including, but not limited to, the bpied
// 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 <string>
#include <iostream>
//#define SHOW_TIME
#ifdef SHOW_TIME
#include <ctime>
#define F(x) x
#else
#define F(x)
#endif
using namespace cv;
using namespace std;
struct CV_GpuMatchTemplateTest: cvtest::BaseTest
{
CV_GpuMatchTemplateTest() {}
void run(int)
{
bool double_ok = gpu::TargetArchs::builtWith(gpu::NATIVE_DOUBLE) &&
gpu::DeviceInfo().supports(gpu::NATIVE_DOUBLE);
if (!double_ok)
{
// For sqrIntegral
ts->printf(cvtest::TS::CONSOLE, "\nCode and device double support is required (CC >= 1.3)");
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
Mat image, templ;
Mat dst_gold;
gpu::GpuMat dst;
int n, m, h, w;
F(clock_t t;)
RNG& rng = ts->get_rng();
for (int cn = 1; cn <= 4; ++cn)
{
F(ts->printf(cvtest::TS::CONSOLE, "cn: %d\n", cn);)
for (int i = 0; i <= 0; ++i)
{
n = rng.uniform(30, 100);
m = rng.uniform(30, 100);
h = rng.uniform(5, n - 1);
w = rng.uniform(5, m - 1);
gen(image, n, m, CV_8U, cn);
gen(templ, h, w, CV_8U, cn);
F(t = clock();)
matchTemplate(image, templ, dst_gold, CV_TM_SQDIFF);
F(cout << "depth: 8U cn: " << cn << " n: " << n << " m: " << m << " w: " << w << " h: " << h << endl;)
F(cout << "cpu:" << clock() - t << endl;)
F(t = clock();)
gpu::matchTemplate(gpu::GpuMat(image), gpu::GpuMat(templ), dst, CV_TM_SQDIFF);
F(cout << "gpu_block: " << clock() - t << endl;)
if (!check(dst_gold, Mat(dst), 5 * h * w * 1e-4f, "SQDIFF 8U")) return;
gen(image, n, m, CV_8U, cn);
gen(templ, h, w, CV_8U, cn);
F(t = clock();)
matchTemplate(image, templ, dst_gold, CV_TM_SQDIFF_NORMED);
F(cout << "depth: 8U cn: " << cn << " n: " << n << " m: " << m << " w: " << w << " h: " << h << endl;)
F(cout << "cpu:" << clock() - t << endl;)
F(t = clock();)
gpu::matchTemplate(gpu::GpuMat(image), gpu::GpuMat(templ), dst, CV_TM_SQDIFF_NORMED);
F(cout << "gpu_block: " << clock() - t << endl;)
if (!check(dst_gold, Mat(dst), h * w * 1e-5f, "SQDIFF_NOREMD 8U")) return;
gen(image, n, m, CV_8U, cn);
gen(templ, h, w, CV_8U, cn);
F(t = clock();)
matchTemplate(image, templ, dst_gold, CV_TM_CCORR);
F(cout << "depth: 8U cn: " << cn << " n: " << n << " m: " << m << " w: " << w << " h: " << h << endl;)
F(cout << "cpu:" << clock() - t << endl;)
F(t = clock();)
gpu::matchTemplate(gpu::GpuMat(image), gpu::GpuMat(templ), dst, CV_TM_CCORR);
F(cout << "gpu_block: " << clock() - t << endl;)
if (!check(dst_gold, Mat(dst), 5 * h * w * cn * cn * 1e-5f, "CCORR 8U")) return;
gen(image, n, m, CV_8U, cn);
gen(templ, h, w, CV_8U, cn);
F(t = clock();)
matchTemplate(image, templ, dst_gold, CV_TM_CCORR_NORMED);
F(cout << "depth: 8U cn: " << cn << " n: " << n << " m: " << m << " w: " << w << " h: " << h << endl;)
F(cout << "cpu:" << clock() - t << endl;)
F(t = clock();)
gpu::matchTemplate(gpu::GpuMat(image), gpu::GpuMat(templ), dst, CV_TM_CCORR_NORMED);
F(cout << "gpu_block: " << clock() - t << endl;)
if (!check(dst_gold, Mat(dst), h * w * 1e-6f, "CCORR_NORMED 8U")) return;
gen(image, n, m, CV_8U, cn);
gen(templ, h, w, CV_8U, cn);
F(t = clock();)
matchTemplate(image, templ, dst_gold, CV_TM_CCOEFF);
F(cout << "depth: 8U cn: " << cn << " n: " << n << " m: " << m << " w: " << w << " h: " << h << endl;)
F(cout << "cpu:" << clock() - t << endl;)
F(t = clock();)
gpu::matchTemplate(gpu::GpuMat(image), gpu::GpuMat(templ), dst, CV_TM_CCOEFF);
F(cout << "gpu_block: " << clock() - t << endl;)
if (!check(dst_gold, Mat(dst), 5 * h * w * cn * cn * 1e-5f, "CCOEFF 8U")) return;
gen(image, n, m, CV_8U, cn);
gen(templ, h, w, CV_8U, cn);
F(t = clock();)
matchTemplate(image, templ, dst_gold, CV_TM_CCOEFF_NORMED);
F(cout << "depth: 8U cn: " << cn << " n: " << n << " m: " << m << " w: " << w << " h: " << h << endl;)
F(cout << "cpu:" << clock() - t << endl;)
F(t = clock();)
gpu::matchTemplate(gpu::GpuMat(image), gpu::GpuMat(templ), dst, CV_TM_CCOEFF_NORMED);
F(cout << "gpu_block: " << clock() - t << endl;)
if (!check(dst_gold, Mat(dst), h * w * 1e-6f, "CCOEFF_NORMED 8U")) return;
gen(image, n, m, CV_32F, cn);
gen(templ, h, w, CV_32F, cn);
F(t = clock();)
matchTemplate(image, templ, dst_gold, CV_TM_SQDIFF);
F(cout << "depth: 32F cn: " << cn << " n: " << n << " m: " << m << " w: " << w << " h: " << h << endl;)
F(cout << "cpu:" << clock() - t << endl;)
F(t = clock();)
gpu::matchTemplate(gpu::GpuMat(image), gpu::GpuMat(templ), dst, CV_TM_SQDIFF);
F(cout << "gpu_block: " << clock() - t << endl;)
if (!check(dst_gold, Mat(dst), 0.25f * h * w * 1e-5f, "SQDIFF 32F")) return;
gen(image, n, m, CV_32F, cn);
gen(templ, h, w, CV_32F, cn);
F(t = clock();)
matchTemplate(image, templ, dst_gold, CV_TM_CCORR);
F(cout << "depth: 32F cn: " << cn << " n: " << n << " m: " << m << " w: " << w << " h: " << h << endl;)
F(cout << "cpu:" << clock() - t << endl;)
F(t = clock();)
gpu::matchTemplate(gpu::GpuMat(image), gpu::GpuMat(templ), dst, CV_TM_CCORR);
F(cout << "gpu_block: " << clock() - t << endl;)
if (!check(dst_gold, Mat(dst), 0.25f * h * w * 1e-5f, "CCORR 32F")) return;
}
}
}
void gen(Mat& a, int rows, int cols, int depth, int cn)
{
RNG rng;
a.create(rows, cols, CV_MAKETYPE(depth, cn));
if (depth == CV_8U)
rng.fill(a, RNG::UNIFORM, Scalar::all(1), Scalar::all(10));
else if (depth == CV_32F)
rng.fill(a, RNG::UNIFORM, Scalar::all(0.001f), Scalar::all(1.f));
}
bool check(const Mat& a, const Mat& b, float max_err, const string& method="")
{
if (a.size() != b.size())
{
ts->printf(cvtest::TS::CONSOLE, "bad size, method=%s\n", method.c_str());
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
//for (int i = 0; i < a.rows; ++i)
//{
// for (int j = 0; j < a.cols; ++j)
// {
// float a_ = a.at<float>(i, j);
// float b_ = b.at<float>(i, j);
// if (fabs(a_ - b_) > max_err)
// {
// ts->printf(cvtest::TS::CONSOLE, "a=%f, b=%f, i=%d, j=%d\n", a_, b_, i, j);
// cin.get();
// }
// }
//}
float err = (float)norm(a, b, NORM_INF);
if (err > max_err)
{
ts->printf(cvtest::TS::CONSOLE, "bad accuracy: %f, method=%s\n", err, method.c_str());
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return false;
}
return true;
}
};
TEST(matchTemplate, accuracy) { CV_GpuMatchTemplateTest test; test.safe_run(); }
struct CV_GpuMatchTemplateFindPatternInBlackTest: cvtest::BaseTest
{
CV_GpuMatchTemplateFindPatternInBlackTest() {}
void run(int)
{
bool double_ok = gpu::TargetArchs::builtWith(gpu::NATIVE_DOUBLE) &&
gpu::DeviceInfo().supports(gpu::NATIVE_DOUBLE);
if (!double_ok)
{
// For sqrIntegral
ts->printf(cvtest::TS::CONSOLE, "\nCode and device double support is required (CC >= 1.3)");
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
Mat image = imread(std::string(ts->get_data_path()) + "matchtemplate/black.png");
if (image.empty())
{
ts->printf(cvtest::TS::CONSOLE, "can't open file '%s'", (std::string(ts->get_data_path())
+ "matchtemplate/black.png").c_str());
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
Mat pattern = imread(std::string(ts->get_data_path()) + "matchtemplate/cat.png");
if (pattern.empty())
{
ts->printf(cvtest::TS::CONSOLE, "can't open file '%s'", (std::string(ts->get_data_path())
+ "matchtemplate/cat.png").c_str());
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
gpu::GpuMat d_image(image);
gpu::GpuMat d_pattern(pattern);
gpu::GpuMat d_result;
double maxValue;
Point maxLoc;
Point maxLocGold(284, 12);
gpu::matchTemplate(d_image, d_pattern, d_result, CV_TM_CCOEFF_NORMED);
gpu::minMaxLoc(d_result, NULL, &maxValue, NULL, &maxLoc );
if (maxLoc != maxLocGold)
{
ts->printf(cvtest::TS::CONSOLE, "bad match (CV_TM_CCOEFF_NORMED): %d %d, must be at: %d %d",
maxLoc.x, maxLoc.y, maxLocGold.x, maxLocGold.y);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
gpu::matchTemplate(d_image, d_pattern, d_result, CV_TM_CCORR_NORMED);
gpu::minMaxLoc(d_result, NULL, &maxValue, NULL, &maxLoc );
if (maxLoc != maxLocGold)
{
ts->printf(cvtest::TS::CONSOLE, "bad match (CV_TM_CCORR_NORMED): %d %d, must be at: %d %d",
maxLoc.x, maxLoc.y, maxLocGold.x, maxLocGold.y);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
}
};
TEST(matchTemplate, find_pattern_in_black) { CV_GpuMatchTemplateFindPatternInBlackTest test; test.safe_run(); }

614
modules/gpu/test/test_matop.cpp Normal file
View File

@ -0,0 +1,614 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other GpuMaterials 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 bpied warranties, including, but not limited to, the bpied
// warranties of merchantability and fitness for a particular purpose are disclaimed.
// In no event shall the Intel Corporation or contributors be liable for any direct,
// indirect, incidental, special, exemplary, or consequential damages
// (including, but not limited to, procurement of substitute goods or services;
// loss of use, data, or profits; or business interruption) however caused
// and on any theory of liability, whether in contract, strict liability,
// or tort (including negligence or otherwise) arising in any way out of
// the use of this software, even if advised of the possibility of such damage.
//
//M*/
#include "test_precomp.hpp"
#ifdef HAVE_CUDA
////////////////////////////////////////////////////////////////////////////////
// merge
struct Merge : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int> >
{
cv::gpu::DeviceInfo devInfo;
int type;
cv::Size size;
std::vector<cv::Mat> src;
cv::Mat dst_gold;
virtual void SetUp()
{
devInfo = std::tr1::get<0>(GetParam());
type = std::tr1::get<1>(GetParam());
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
size = cv::Size(rng.uniform(20, 150), rng.uniform(20, 150));
int depth = CV_MAT_DEPTH(type);
int num_channels = CV_MAT_CN(type);
src.reserve(num_channels);
for (int i = 0; i < num_channels; ++i)
src.push_back(cv::Mat(size, depth, cv::Scalar::all(i)));
cv::merge(src, dst_gold);
}
};
TEST_P(Merge, Accuracy)
{
if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(type);
PRINT_PARAM(size);
cv::Mat dst;
ASSERT_NO_THROW(
std::vector<cv::gpu::GpuMat> dev_src;
cv::gpu::GpuMat dev_dst;
for (size_t i = 0; i < src.size(); ++i)
dev_src.push_back(cv::gpu::GpuMat(src[i]));
cv::gpu::merge(dev_src, dev_dst);
dev_dst.download(dst);
);
EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
}
INSTANTIATE_TEST_CASE_P(MatOp, Merge, testing::Combine(
testing::ValuesIn(devices()),
testing::ValuesIn(all_types())));
////////////////////////////////////////////////////////////////////////////////
// split
struct Split : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int> >
{
cv::gpu::DeviceInfo devInfo;
int type;
cv::Size size;
cv::Mat src;
std::vector<cv::Mat> dst_gold;
virtual void SetUp()
{
devInfo = std::tr1::get<0>(GetParam());
type = std::tr1::get<1>(GetParam());
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
size = cv::Size(rng.uniform(20, 150), rng.uniform(20, 150));
src.create(size, type);
src.setTo(cv::Scalar(1.0, 2.0, 3.0, 4.0));
cv::split(src, dst_gold);
}
};
TEST_P(Split, Accuracy)
{
if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(type);
PRINT_PARAM(size);
std::vector<cv::Mat> dst;
ASSERT_NO_THROW(
std::vector<cv::gpu::GpuMat> dev_dst;
cv::gpu::split(cv::gpu::GpuMat(src), dev_dst);
dst.resize(dev_dst.size());
for (size_t i = 0; i < dev_dst.size(); ++i)
dev_dst[i].download(dst[i]);
);
ASSERT_EQ(dst_gold.size(), dst.size());
for (size_t i = 0; i < dst_gold.size(); ++i)
{
EXPECT_MAT_NEAR(dst_gold[i], dst[i], 0.0);
}
}
INSTANTIATE_TEST_CASE_P(MatOp, Split, testing::Combine(
testing::ValuesIn(devices()),
testing::ValuesIn(all_types())));
////////////////////////////////////////////////////////////////////////////////
// split_merge_consistency
struct SplitMerge : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int> >
{
cv::gpu::DeviceInfo devInfo;
int type;
cv::Size size;
cv::Mat orig;
virtual void SetUp()
{
devInfo = std::tr1::get<0>(GetParam());
type = std::tr1::get<1>(GetParam());
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
size = cv::Size(rng.uniform(20, 150), rng.uniform(20, 150));
orig.create(size, type);
orig.setTo(cv::Scalar(1.0, 2.0, 3.0, 4.0));
}
};
TEST_P(SplitMerge, Consistency)
{
if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(type);
PRINT_PARAM(size);
cv::Mat final;
ASSERT_NO_THROW(
std::vector<cv::gpu::GpuMat> dev_vec;
cv::gpu::GpuMat dev_final;
cv::gpu::split(cv::gpu::GpuMat(orig), dev_vec);
cv::gpu::merge(dev_vec, dev_final);
dev_final.download(final);
);
EXPECT_MAT_NEAR(orig, final, 0.0);
}
INSTANTIATE_TEST_CASE_P(MatOp, SplitMerge, testing::Combine(
testing::ValuesIn(devices()),
testing::ValuesIn(all_types())));
////////////////////////////////////////////////////////////////////////////////
// setTo
struct SetTo : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int> >
{
cv::gpu::DeviceInfo devInfo;
int type;
cv::Size size;
cv::Mat mat_gold;
virtual void SetUp()
{
devInfo = std::tr1::get<0>(GetParam());
type = std::tr1::get<1>(GetParam());
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
size = cv::Size(rng.uniform(20, 150), rng.uniform(20, 150));
mat_gold.create(size, type);
}
};
TEST_P(SetTo, Zero)
{
if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(type);
PRINT_PARAM(size);
static cv::Scalar zero = cv::Scalar::all(0);
cv::Mat mat;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_mat(mat_gold);
mat_gold.setTo(zero);
dev_mat.setTo(zero);
dev_mat.download(mat);
);
EXPECT_MAT_NEAR(mat_gold, mat, 0.0);
}
TEST_P(SetTo, SameVal)
{
if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(type);
PRINT_PARAM(size);
static cv::Scalar s = cv::Scalar::all(1);
cv::Mat mat;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_mat(mat_gold);
mat_gold.setTo(s);
dev_mat.setTo(s);
dev_mat.download(mat);
);
EXPECT_MAT_NEAR(mat_gold, mat, 0.0);
}
TEST_P(SetTo, DifferentVal)
{
if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(type);
PRINT_PARAM(size);
static cv::Scalar s = cv::Scalar(1, 2, 3, 4);
cv::Mat mat;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_mat(mat_gold);
mat_gold.setTo(s);
dev_mat.setTo(s);
dev_mat.download(mat);
);
EXPECT_MAT_NEAR(mat_gold, mat, 0.0);
}
TEST_P(SetTo, Masked)
{
if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(type);
PRINT_PARAM(size);
static cv::Scalar s = cv::Scalar(1, 2, 3, 4);
cv::Mat mat;
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
cv::Mat mask = cvtest::randomMat(rng, mat.size(), CV_8UC1, 0.0, 1.5, false);
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_mat(mat_gold);
mat_gold.setTo(s, mask);
dev_mat.setTo(s, cv::gpu::GpuMat(mask));
dev_mat.download(mat);
);
EXPECT_MAT_NEAR(mat_gold, mat, 0.0);
}
INSTANTIATE_TEST_CASE_P(MatOp, SetTo, testing::Combine(
testing::ValuesIn(devices()),
testing::ValuesIn(all_types())));
////////////////////////////////////////////////////////////////////////////////
// copyTo
struct CopyTo : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int> >
{
cv::gpu::DeviceInfo devInfo;
int type;
cv::Size size;
cv::Mat src;
virtual void SetUp()
{
devInfo = std::tr1::get<0>(GetParam());
type = std::tr1::get<1>(GetParam());
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
size = cv::Size(rng.uniform(20, 150), rng.uniform(20, 150));
src = cvtest::randomMat(rng, size, type, 0.0, 127.0, false);
}
};
TEST_P(CopyTo, WithoutMask)
{
if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(type);
PRINT_PARAM(size);
cv::Mat dst_gold;
src.copyTo(dst_gold);
cv::Mat dst;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_src(src);
cv::gpu::GpuMat dev_dst;
dev_src.copyTo(dev_dst);
dev_dst.download(dst);
);
EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
}
TEST_P(CopyTo, Masked)
{
if (CV_MAT_DEPTH(type) == CV_64F && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(type);
PRINT_PARAM(size);
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
cv::Mat mask = cvtest::randomMat(rng, src.size(), CV_8UC1, 0.0, 1.5, false);
cv::Mat dst_gold;
src.copyTo(dst_gold, mask);
cv::Mat dst;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_src(src);
cv::gpu::GpuMat dev_dst;
dev_src.copyTo(dev_dst, cv::gpu::GpuMat(mask));
dev_dst.download(dst);
);
EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
}
INSTANTIATE_TEST_CASE_P(MatOp, CopyTo, testing::Combine(
testing::ValuesIn(devices()),
testing::ValuesIn(all_types())));
////////////////////////////////////////////////////////////////////////////////
// convertTo
struct ConvertTo : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int, int> >
{
cv::gpu::DeviceInfo devInfo;
int depth1;
int depth2;
cv::Size size;
cv::Mat src;
virtual void SetUp()
{
devInfo = std::tr1::get<0>(GetParam());
depth1 = std::tr1::get<1>(GetParam());
depth2 = std::tr1::get<2>(GetParam());
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
size = cv::Size(rng.uniform(20, 150), rng.uniform(20, 150));
src = cvtest::randomMat(rng, size, depth1, 0.0, 127.0, false);
}
};
TEST_P(ConvertTo, WithoutScaling)
{
if ((depth1 == CV_64F || depth2 == CV_64F) && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(depth1);
PRINT_TYPE(depth2);
PRINT_PARAM(size);
cv::Mat dst_gold;
src.convertTo(dst_gold, depth2);
cv::Mat dst;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_src(src);
cv::gpu::GpuMat dev_dst;
dev_src.convertTo(dev_dst, depth2);
dev_dst.download(dst);
);
EXPECT_MAT_NEAR(dst_gold, dst, 0.0);
}
TEST_P(ConvertTo, WithScaling)
{
if ((depth1 == CV_64F || depth2 == CV_64F) && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
return;
PRINT_PARAM(devInfo);
PRINT_TYPE(depth1);
PRINT_TYPE(depth2);
PRINT_PARAM(size);
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
const double a = rng.uniform(0.0, 1.0);
const double b = rng.uniform(-10.0, 10.0);
PRINT_PARAM(a);
PRINT_PARAM(b);
cv::Mat dst_gold;
src.convertTo(dst_gold, depth2, a, b);
cv::Mat dst;
ASSERT_NO_THROW(
cv::gpu::GpuMat dev_src(src);
cv::gpu::GpuMat dev_dst;
dev_src.convertTo(dev_dst, depth2, a, b);
dev_dst.download(dst);
);
const double eps = depth2 < CV_32F ? 1 : 1e-4;
EXPECT_MAT_NEAR(dst_gold, dst, eps);
}
INSTANTIATE_TEST_CASE_P(MatOp, ConvertTo, testing::Combine(
testing::ValuesIn(devices()),
testing::ValuesIn(types(CV_8U, CV_64F, 1, 1)),
testing::ValuesIn(types(CV_8U, CV_64F, 1, 1))));
////////////////////////////////////////////////////////////////////////////////
// async
struct Async : testing::TestWithParam<cv::gpu::DeviceInfo>
{
cv::gpu::DeviceInfo devInfo;
cv::gpu::CudaMem src;
cv::Mat dst_gold0;
cv::Mat dst_gold1;
virtual void SetUp()
{
devInfo = GetParam();
cv::gpu::setDevice(devInfo.deviceID());
cv::RNG& rng = cvtest::TS::ptr()->get_rng();
int rows = rng.uniform(100, 200);
int cols = rng.uniform(100, 200);
src = cv::gpu::CudaMem(cv::Mat::zeros(rows, cols, CV_8UC1));
dst_gold0 = cv::Mat(rows, cols, CV_8UC1, cv::Scalar::all(255));
dst_gold1 = cv::Mat(rows, cols, CV_8UC1, cv::Scalar::all(128));
}
};
TEST_P(Async, Accuracy)
{
PRINT_PARAM(devInfo);
cv::Mat dst0, dst1;
ASSERT_NO_THROW(
cv::gpu::CudaMem cpudst0;
cv::gpu::CudaMem cpudst1;
cv::gpu::GpuMat gpusrc;
cv::gpu::GpuMat gpudst0;
cv::gpu::GpuMat gpudst1(src.rows, src.cols, CV_8UC1);
cv::gpu::Stream stream0;
cv::gpu::Stream stream1;
stream0.enqueueUpload(src, gpusrc);
cv::gpu::bitwise_not(gpusrc, gpudst0, cv::gpu::GpuMat(), stream0);
stream0.enqueueDownload(gpudst0, cpudst0);
stream1.enqueueMemSet(gpudst1, cv::Scalar::all(128));
stream1.enqueueDownload(gpudst1, cpudst1);
stream0.waitForCompletion();
stream1.waitForCompletion();
dst0 = cpudst0.createMatHeader();
dst1 = cpudst1.createMatHeader();
);
EXPECT_MAT_NEAR(dst_gold0, dst0, 0.0);
EXPECT_MAT_NEAR(dst_gold1, dst1, 0.0);
}
INSTANTIATE_TEST_CASE_P(MatOp, Async, testing::ValuesIn(devices()));
#endif // HAVE_CUDA

233
modules/gpu/test/test_meanshift.cpp
View File

@ -1,233 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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 <iostream>
#include <string>
using namespace cv;
using namespace cv::gpu;
struct CV_GpuMeanShiftTest : public cvtest::BaseTest
{
CV_GpuMeanShiftTest() {}
void run(int)
{
bool cc12_ok = TargetArchs::builtWith(FEATURE_SET_COMPUTE_12) && DeviceInfo().supports(FEATURE_SET_COMPUTE_12);
if (!cc12_ok)
{
ts->printf(cvtest::TS::CONSOLE, "\nCompute capability 1.2 is required");
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
int spatialRad = 30;
int colorRad = 30;
cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "meanshift/cones.png");
cv::Mat img_template;
if (cv::gpu::TargetArchs::builtWith(cv::gpu::FEATURE_SET_COMPUTE_20) &&
cv::gpu::DeviceInfo().supports(cv::gpu::FEATURE_SET_COMPUTE_20))
img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result.png");
else
img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result_CC1X.png");
if (img.empty() || img_template.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
cv::Mat rgba;
cvtColor(img, rgba, CV_BGR2BGRA);
cv::gpu::GpuMat res;
cv::gpu::meanShiftFiltering( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad );
if (res.type() != CV_8UC4)
{
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
cv::Mat result;
res.download(result);
uchar maxDiff = 0;
for (int j = 0; j < result.rows; ++j)
{
const uchar* res_line = result.ptr<uchar>(j);
const uchar* ref_line = img_template.ptr<uchar>(j);
for (int i = 0; i < result.cols; ++i)
{
for (int k = 0; k < 3; ++k)
{
const uchar& ch1 = res_line[result.channels()*i + k];
const uchar& ch2 = ref_line[img_template.channels()*i + k];
uchar diff = static_cast<uchar>(abs(ch1 - ch2));
if (maxDiff < diff)
maxDiff = diff;
}
}
}
if (maxDiff > 0)
{
ts->printf(cvtest::TS::LOG, "\nMeanShift maxDiff = %d\n", maxDiff);
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
ts->set_failed_test_info(cvtest::TS::OK);
}
};
TEST(meanShift, accuracy) { CV_GpuMeanShiftTest test; test.safe_run(); }
struct CV_GpuMeanShiftProcTest : public cvtest::BaseTest
{
CV_GpuMeanShiftProcTest() {}
void run(int)
{
bool cc12_ok = TargetArchs::builtWith(FEATURE_SET_COMPUTE_12) && DeviceInfo().supports(FEATURE_SET_COMPUTE_12);
if (!cc12_ok)
{
ts->printf(cvtest::TS::CONSOLE, "\nCompute capability 1.2 is required");
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
int spatialRad = 30;
int colorRad = 30;
cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "meanshift/cones.png");
if (img.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
cv::Mat rgba;
cvtColor(img, rgba, CV_BGR2BGRA);
cv::gpu::GpuMat h_rmap_filtered;
cv::gpu::meanShiftFiltering( cv::gpu::GpuMat(rgba), h_rmap_filtered, spatialRad, colorRad );
cv::gpu::GpuMat d_rmap;
cv::gpu::GpuMat d_spmap;
cv::gpu::meanShiftProc( cv::gpu::GpuMat(rgba), d_rmap, d_spmap, spatialRad, colorRad );
if (d_rmap.type() != CV_8UC4)
{
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
cv::Mat rmap_filtered;
h_rmap_filtered.download(rmap_filtered);
cv::Mat rmap;
d_rmap.download(rmap);
uchar maxDiff = 0;
for (int j = 0; j < rmap_filtered.rows; ++j)
{
const uchar* res_line = rmap_filtered.ptr<uchar>(j);
const uchar* ref_line = rmap.ptr<uchar>(j);
for (int i = 0; i < rmap_filtered.cols; ++i)
{
for (int k = 0; k < 3; ++k)
{
const uchar& ch1 = res_line[rmap_filtered.channels()*i + k];
const uchar& ch2 = ref_line[rmap.channels()*i + k];
uchar diff = static_cast<uchar>(abs(ch1 - ch2));
if (maxDiff < diff)
maxDiff = diff;
}
}
}
if (maxDiff > 0)
{
ts->printf(cvtest::TS::LOG, "\nMeanShiftProc maxDiff = %d\n", maxDiff);
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
cv::Mat spmap;
d_spmap.download(spmap);
cv::Mat spmap_template;
cv::FileStorage fs;
if (cv::gpu::TargetArchs::builtWith(cv::gpu::FEATURE_SET_COMPUTE_20) &&
cv::gpu::DeviceInfo().supports(cv::gpu::FEATURE_SET_COMPUTE_20))
fs.open(std::string(ts->get_data_path()) + "meanshift/spmap.yaml", cv::FileStorage::READ);
else
fs.open(std::string(ts->get_data_path()) + "meanshift/spmap_CC1X.yaml", cv::FileStorage::READ);
fs["spmap"] >> spmap_template;
for (int y = 0; y < spmap.rows; ++y) {
for (int x = 0; x < spmap.cols; ++x) {
cv::Point_<short> expected = spmap_template.at<cv::Point_<short> >(y, x);
cv::Point_<short> actual = spmap.at<cv::Point_<short> >(y, x);
int diff = (expected - actual).dot(expected - actual);
if (actual != expected) {
ts->printf(cvtest::TS::LOG, "\nMeanShiftProc SpMap is bad, diff=%d\n", diff);
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
}
}
ts->set_failed_test_info(cvtest::TS::OK);
}
};
TEST(meanShiftProc, accuracy) { CV_GpuMeanShiftProcTest test; test.safe_run(); }

122
modules/gpu/test/test_mssegmentation.cpp
View File

@ -1,122 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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 <iostream>
#include <string>
#include <iosfwd>
#include "test_precomp.hpp"
using namespace cv;
using namespace cv::gpu;
using namespace std;
struct CV_GpuMeanShiftSegmentationTest : public cvtest::BaseTest {
CV_GpuMeanShiftSegmentationTest() {}
void run(int)
{
bool cc12_ok = TargetArchs::builtWith(FEATURE_SET_COMPUTE_12) && DeviceInfo().supports(FEATURE_SET_COMPUTE_12);
if (!cc12_ok)
{
ts->printf(cvtest::TS::CONSOLE, "\nCompute capability 1.2 is required");
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
Mat img_rgb = imread(string(ts->get_data_path()) + "meanshift/cones.png");
if (img_rgb.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
Mat img;
cvtColor(img_rgb, img, CV_BGR2BGRA);
for (int minsize = 0; minsize < 2000; minsize = (minsize + 1) * 4)
{
stringstream path;
path << ts->get_data_path() << "meanshift/cones_segmented_sp10_sr10_minsize" << minsize;
if (TargetArchs::builtWith(FEATURE_SET_COMPUTE_20) && DeviceInfo().supports(FEATURE_SET_COMPUTE_20))
path << ".png";
else
path << "_CC1X.png";
Mat dst;
meanShiftSegmentation((GpuMat)img, dst, 10, 10, minsize);
Mat dst_rgb;
cvtColor(dst, dst_rgb, CV_BGRA2BGR);
//imwrite(path.str(), dst_rgb);
Mat dst_ref = imread(path.str());
if (dst_ref.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
if (CheckSimilarity(dst_rgb, dst_ref, 1e-3f) != cvtest::TS::OK)
{
ts->printf(cvtest::TS::LOG, "\ndiffers from image *minsize%d.png\n", minsize);
ts->set_failed_test_info(cvtest::TS::FAIL_BAD_ACCURACY);
}
}
ts->set_failed_test_info(cvtest::TS::OK);
}
int CheckSimilarity(const Mat& m1, const Mat& m2, float max_err)
{
Mat diff;
cv::matchTemplate(m1, m2, diff, CV_TM_CCORR_NORMED);
float err = abs(diff.at<float>(0, 0) - 1.f);
if (err > max_err)
return cvtest::TS::FAIL_INVALID_OUTPUT;
return cvtest::TS::OK;
}
};
TEST(meanShiftSegmentation, regression) { CV_GpuMeanShiftSegmentationTest test; test.safe_run(); }

219
modules/gpu/test/test_nvidia.cpp
View File

@ -40,35 +40,198 @@
//M*/
#include "test_precomp.hpp"
#include "cvconfig.h"
class CV_NVidiaTestsCaller : public cvtest::BaseTest
#ifdef HAVE_CUDA
enum OutputLevel
{
public:
CV_NVidiaTestsCaller() {}
virtual ~CV_NVidiaTestsCaller() {}
protected:
void run( int )
{
;
#if defined(HAVE_CUDA)
bool main_nvidia(const std::string&);
// Invoke all NVIDIA Staging tests and obtain the result
bool passed = main_nvidia(std::string(ts->get_data_path()) + "haarcascade/");
if (passed)
ts->set_failed_test_info(cvtest::TS::OK);
else
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
#else
ts->set_failed_test_info(cvtest::TS::SKIPPED);
#endif
}
OutputLevelNone,
OutputLevelCompact,
OutputLevelFull
};
TEST(NVidia, multitest) { CV_NVidiaTestsCaller test; test.safe_run(); }
bool nvidia_NPPST_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel);
bool nvidia_NPPST_Squared_Integral_Image(const std::string& test_data_path, OutputLevel outputLevel);
bool nvidia_NPPST_RectStdDev(const std::string& test_data_path, OutputLevel outputLevel);
bool nvidia_NPPST_Resize(const std::string& test_data_path, OutputLevel outputLevel);
bool nvidia_NPPST_Vector_Operations(const std::string& test_data_path, OutputLevel outputLevel);
bool nvidia_NPPST_Transpose(const std::string& test_data_path, OutputLevel outputLevel);
bool nvidia_NCV_Vector_Operations(const std::string& test_data_path, OutputLevel outputLevel);
bool nvidia_NCV_Haar_Cascade_Loader(const std::string& test_data_path, OutputLevel outputLevel);
bool nvidia_NCV_Haar_Cascade_Application(const std::string& test_data_path, OutputLevel outputLevel);
bool nvidia_NCV_Hypotheses_Filtration(const std::string& test_data_path, OutputLevel outputLevel);
bool nvidia_NCV_Visualization(const std::string& test_data_path, OutputLevel outputLevel);
struct NVidiaTest : testing::TestWithParam<cv::gpu::DeviceInfo>
{
static std::string path;
cv::gpu::DeviceInfo devInfo;
static void SetUpTestCase()
{
path = std::string(cvtest::TS::ptr()->get_data_path()) + "haarcascade/";
}
virtual void SetUp()
{
devInfo = GetParam();
cv::gpu::setDevice(devInfo.deviceID());
}
};
std::string NVidiaTest::path;
struct NPPST : NVidiaTest {};
struct NCV : NVidiaTest {};
OutputLevel nvidiaTestOutputLevel = OutputLevelNone;
TEST_P(NPPST, Integral)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NPPST_Integral_Image(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
TEST_P(NPPST, SquaredIntegral)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NPPST_Squared_Integral_Image(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
TEST_P(NPPST, RectStdDev)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NPPST_RectStdDev(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
TEST_P(NPPST, Resize)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NPPST_Resize(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
TEST_P(NPPST, VectorOperations)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NPPST_Vector_Operations(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
TEST_P(NPPST, Transpose)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NPPST_Transpose(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
TEST_P(NCV, VectorOperations)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NCV_Vector_Operations(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
TEST_P(NCV, HaarCascadeLoader)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NCV_Haar_Cascade_Loader(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
TEST_P(NCV, HaarCascadeApplication)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NCV_Haar_Cascade_Application(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
TEST_P(NCV, HypothesesFiltration)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NCV_Hypotheses_Filtration(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
TEST_P(NCV, Visualization)
{
PRINT_PARAM(devInfo);
bool res;
ASSERT_NO_THROW(
res = nvidia_NCV_Visualization(path, nvidiaTestOutputLevel);
);
ASSERT_TRUE(res);
}
INSTANTIATE_TEST_CASE_P(NVidia, NPPST, testing::ValuesIn(devices()));
INSTANTIATE_TEST_CASE_P(NVidia, NCV, testing::ValuesIn(devices()));
#endif // HAVE_CUDA

84
modules/gpu/test/test_operator_async_call.cpp
View File

@ -1,84 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
using namespace std;
using namespace cv;
using namespace cv::gpu;
struct CV_AsyncGpuMatTest : public cvtest::BaseTest
{
CV_AsyncGpuMatTest() {}
void run(int)
{
CudaMem src(Mat::zeros(100, 100, CV_8UC1));
GpuMat gpusrc;
GpuMat gpudst0, gpudst1(100, 100, CV_8UC1);
CudaMem cpudst0;
CudaMem cpudst1;
Stream stream0, stream1;
stream0.enqueueUpload(src, gpusrc);
bitwise_not(gpusrc, gpudst0, GpuMat(), stream0);
stream0.enqueueDownload(gpudst0, cpudst0);
stream1.enqueueMemSet(gpudst1, Scalar::all(128));
stream1.enqueueDownload(gpudst1, cpudst1);
stream0.waitForCompletion();
stream1.waitForCompletion();
Mat cpu_gold0(100, 100, CV_8UC1, Scalar::all(255));
Mat cpu_gold1(100, 100, CV_8UC1, Scalar::all(128));
if (norm((Mat)cpudst0, cpu_gold0, NORM_INF) > 0 || norm((Mat)cpudst1, cpu_gold1, NORM_INF) > 0)
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
else
ts->set_failed_test_info(cvtest::TS::OK);
}
};
TEST(GpuMat, async) { CV_AsyncGpuMatTest test; test.safe_run(); }

110
modules/gpu/test/test_operator_convert_to.cpp
View File

@ -1,110 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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 <fstream>
#include <iterator>
#include <numeric>
using namespace cv;
using namespace std;
using namespace gpu;
class CV_GpuMatOpConvertToTest : public cvtest::BaseTest
{
public:
CV_GpuMatOpConvertToTest() {}
~CV_GpuMatOpConvertToTest() {}
protected:
void run(int);
};
void CV_GpuMatOpConvertToTest::run(int /* start_from */)
{
const Size img_size(67, 35);
const char* types_str[] = {"CV_8U", "CV_8S", "CV_16U", "CV_16S", "CV_32S", "CV_32F", "CV_64F"};
bool passed = true;
int lastType = CV_32F;
if (TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE))
lastType = CV_64F;
for (int i = 0; i <= lastType && passed; ++i)
{
for (int j = 0; j <= lastType && passed; ++j)
{
for (int c = 1; c < 5 && passed; ++c)
{
const int src_type = CV_MAKETYPE(i, c);
const int dst_type = j;
cv::RNG& rng = ts->get_rng();
Mat cpumatsrc(img_size, src_type);
rng.fill(cpumatsrc, RNG::UNIFORM, Scalar::all(0), Scalar::all(300));
GpuMat gpumatsrc(cpumatsrc);
Mat cpumatdst;
GpuMat gpumatdst;
cpumatsrc.convertTo(cpumatdst, dst_type, 0.5, 3.0);
gpumatsrc.convertTo(gpumatdst, dst_type, 0.5, 3.0);
double r = norm(cpumatdst, (Mat)gpumatdst, NORM_INF);
if (r > 1)
{
ts->printf(cvtest::TS::LOG,
"\nFAILED: SRC_TYPE=%sC%d DST_TYPE=%s NORM = %f\n",
types_str[i], c, types_str[j], r);
passed = false;
}
}
}
}
ts->set_failed_test_info(passed ? cvtest::TS::OK : cvtest::TS::FAIL_GENERIC);
}
TEST(GpuMat_convertTo, accuracy) { CV_GpuMatOpConvertToTest test; test.safe_run(); }

145
modules/gpu/test/test_operator_copy_to.cpp
View File

@ -1,145 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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 <fstream>
#include <iterator>
#include <numeric>
#include <iomanip> // for cout << setw()
using namespace cv;
using namespace std;
using namespace gpu;
class CV_GpuMatOpCopyToTest : public cvtest::BaseTest
{
public:
CV_GpuMatOpCopyToTest()
{
rows = 234;
cols = 123;
}
~CV_GpuMatOpCopyToTest() {}
protected:
void run(int);
template <typename T>
void print_mat(const T & mat, const std::string & name) const;
bool compare_matrix(cv::Mat & cpumat, gpu::GpuMat & gpumat);
private:
int rows;
int cols;
};
template<typename T>
void CV_GpuMatOpCopyToTest::print_mat(const T & mat, const std::string & name) const { cv::imshow(name, mat); }
bool CV_GpuMatOpCopyToTest::compare_matrix(cv::Mat & cpumat, gpu::GpuMat & gpumat)
{
Mat cmat(cpumat.size(), cpumat.type(), Scalar::all(0));
GpuMat gmat(cmat);
Mat cpumask(cpumat.size(), CV_8U);
cv::RNG& rng = ts->get_rng();
rng.fill(cpumask, RNG::NORMAL, Scalar::all(0), Scalar::all(127));
threshold(cpumask, cpumask, 0, 127, THRESH_BINARY);
GpuMat gpumask(cpumask);
//int64 time = getTickCount();
cpumat.copyTo(cmat, cpumask);
//int64 time1 = getTickCount();
gpumat.copyTo(gmat, gpumask);
//int64 time2 = getTickCount();
//std::cout << "\ntime cpu: " << std::fixed << std::setprecision(12) << 1.0 / double((time1 - time) / (double)getTickFrequency());
//std::cout << "\ntime gpu: " << std::fixed << std::setprecision(12) << 1.0 / double((time2 - time1) / (double)getTickFrequency());
//std::cout << "\n";
#ifdef PRINT_MATRIX
print_mat(cmat, "cpu mat");
print_mat(gmat, "gpu mat");
print_mat(cpumask, "cpu mask");
print_mat(gpumask, "gpu mask");
cv::waitKey(0);
#endif
double ret = norm(cmat, (Mat)gmat);
if (ret < 1.0)
return true;
else
{
ts->printf(cvtest::TS::LOG, "\nNorm: %f\n", ret);
return false;
}
}
void CV_GpuMatOpCopyToTest::run( int /* start_from */)
{
bool is_test_good = true;
int lastType = CV_32F;
if (TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE))
lastType = CV_64F;
for (int i = 0 ; i <= lastType; i++)
{
Mat cpumat(rows, cols, i);
cpumat.setTo(Scalar::all(127));
GpuMat gpumat(cpumat);
is_test_good &= compare_matrix(cpumat, gpumat);
}
if (is_test_good == true)
ts->set_failed_test_info(cvtest::TS::OK);
else
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
}
TEST(GpuMat_copyTo, accuracy) { CV_GpuMatOpCopyToTest test; test.safe_run(); }

123
modules/gpu/test/test_operator_set_to.cpp
View File

@ -1,123 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
using namespace cv;
using namespace std;
using namespace gpu;
class CV_GpuMatOpSetToTest : public cvtest::BaseTest
{
public:
CV_GpuMatOpSetToTest();
~CV_GpuMatOpSetToTest() {}
protected:
void run(int);
bool testSetTo(cv::Mat& cpumat, gpu::GpuMat& gpumat, const cv::Mat& cpumask = cv::Mat(), const cv::gpu::GpuMat& gpumask = cv::gpu::GpuMat());
private:
int rows;
int cols;
Scalar s;
};
CV_GpuMatOpSetToTest::CV_GpuMatOpSetToTest()
{
rows = 35;
cols = 67;
s.val[0] = 127.0;
s.val[1] = 127.0;
s.val[2] = 127.0;
s.val[3] = 127.0;
}
bool CV_GpuMatOpSetToTest::testSetTo(cv::Mat& cpumat, gpu::GpuMat& gpumat, const cv::Mat& cpumask, const cv::gpu::GpuMat& gpumask)
{
cpumat.setTo(s, cpumask);
gpumat.setTo(s, gpumask);
double ret = norm(cpumat, (Mat)gpumat, NORM_INF);
if (ret < std::numeric_limits<double>::epsilon())
return true;
else
{
ts->printf(cvtest::TS::LOG, "\nNorm: %f\n", ret);
return false;
}
}
void CV_GpuMatOpSetToTest::run( int /* start_from */)
{
bool is_test_good = true;
cv::Mat cpumask(rows, cols, CV_8UC1);
cv::RNG& rng = ts->get_rng();
rng.fill(cpumask, RNG::UNIFORM, cv::Scalar::all(0.0), cv::Scalar(1.5));
cv::gpu::GpuMat gpumask(cpumask);
int lastType = CV_32F;
if (TargetArchs::builtWith(NATIVE_DOUBLE) && DeviceInfo().supports(NATIVE_DOUBLE))
lastType = CV_64F;
for (int i = 0; i <= lastType; i++)
{
for (int cn = 1; cn <= 4; ++cn)
{
int mat_type = CV_MAKETYPE(i, cn);
Mat cpumat(rows, cols, mat_type, Scalar::all(0));
GpuMat gpumat(cpumat);
is_test_good &= testSetTo(cpumat, gpumat, cpumask, gpumask);
}
}
if (is_test_good == true)
ts->set_failed_test_info(cvtest::TS::OK);
else
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
}
TEST(GpuMat_setTo, accuracy) { CV_GpuMatOpSetToTest test; test.safe_run(); }

41
modules/gpu/test/test_precomp.cpp
View File

@ -1 +1,42 @@
/*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.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"

51
modules/gpu/test/test_precomp.hpp
View File

@ -1,13 +1,62 @@
/*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.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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*/
#ifndef __OPENCV_TEST_PRECOMP_HPP__
#define __OPENCV_TEST_PRECOMP_HPP__
#include <iostream>
#include <cmath>
#include <cstdio>
#include <fstream>
#include <sstream>
#include <limits>
#include <string>
#include <algorithm>
#include <iterator>
#include "cvconfig.h"
#include "opencv2/core/core.hpp"
#include "opencv2/highgui/highgui.hpp"
#include "opencv2/calib3d/calib3d.hpp"
#include "opencv2/imgproc/imgproc.hpp"
#include "opencv2/ts/ts.hpp"
#include "opencv2/gpu/gpu.hpp"
#include "test_gpu_base.hpp"
#endif

312
modules/gpu/test/test_split_merge.cpp
View File

@ -1,312 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
// Copyright (C) 2009, Willow Garage Inc., all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other GpuMaterials 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 bpied warranties, including, but not limited to, the bpied
// 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"
using namespace std;
using namespace cv;
////////////////////////////////////////////////////////////////////////////////
// Merge
struct CV_MergeTest : public cvtest::BaseTest
{
void can_merge(size_t rows, size_t cols);
void can_merge_submatrixes(size_t rows, size_t cols);
void run(int);
};
void CV_MergeTest::can_merge(size_t rows, size_t cols)
{
bool double_ok = gpu::TargetArchs::builtWith(gpu::NATIVE_DOUBLE) &&
gpu::DeviceInfo().supports(gpu::NATIVE_DOUBLE);
size_t depth_end = double_ok ? CV_64F : CV_32F;
for (size_t num_channels = 1; num_channels <= 4; ++num_channels)
for (size_t depth = CV_8U; depth <= depth_end; ++depth)
{
vector<Mat> src;
for (size_t i = 0; i < num_channels; ++i)
src.push_back(Mat(rows, cols, depth, Scalar::all(static_cast<double>(i))));
Mat dst(rows, cols, CV_MAKETYPE(depth, num_channels));
cv::merge(src, dst);
vector<gpu::GpuMat> dev_src;
for (size_t i = 0; i < num_channels; ++i)
dev_src.push_back(gpu::GpuMat(src[i]));
gpu::GpuMat dev_dst(rows, cols, CV_MAKETYPE(depth, num_channels));
cv::gpu::merge(dev_src, dev_dst);
Mat host_dst = dev_dst;
double err = norm(dst, host_dst, NORM_INF);
if (err > 1e-3)
{
//ts->printf(cvtest::TS::CONSOLE, "\nNorm: %f\n", err);
//ts->printf(cvtest::TS::CONSOLE, "Depth: %d\n", depth);
//ts->printf(cvtest::TS::CONSOLE, "Rows: %d\n", rows);
//ts->printf(cvtest::TS::CONSOLE, "Cols: %d\n", cols);
//ts->printf(cvtest::TS::CONSOLE, "NumChannels: %d\n", num_channels);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
}
}
void CV_MergeTest::can_merge_submatrixes(size_t rows, size_t cols)
{
bool double_ok = gpu::TargetArchs::builtWith(gpu::NATIVE_DOUBLE) &&
gpu::DeviceInfo().supports(gpu::NATIVE_DOUBLE);
size_t depth_end = double_ok ? CV_64F : CV_32F;
for (size_t num_channels = 1; num_channels <= 4; ++num_channels)
for (size_t depth = CV_8U; depth <= depth_end; ++depth)
{
vector<Mat> src;
for (size_t i = 0; i < num_channels; ++i)
{
Mat m(rows * 2, cols * 2, depth, Scalar::all(static_cast<double>(i)));
src.push_back(m(Range(rows / 2, rows / 2 + rows), Range(cols / 2, cols / 2 + cols)));
}
Mat dst(rows, cols, CV_MAKETYPE(depth, num_channels));
cv::merge(src, dst);
vector<gpu::GpuMat> dev_src;
for (size_t i = 0; i < num_channels; ++i)
dev_src.push_back(gpu::GpuMat(src[i]));
gpu::GpuMat dev_dst(rows, cols, CV_MAKETYPE(depth, num_channels));
cv::gpu::merge(dev_src, dev_dst);
Mat host_dst = dev_dst;
double err = norm(dst, host_dst, NORM_INF);
if (err > 1e-3)
{
//ts->printf(cvtest::TS::CONSOLE, "\nNorm: %f\n", err);
//ts->printf(cvtest::TS::CONSOLE, "Depth: %d\n", depth);
//ts->printf(cvtest::TS::CONSOLE, "Rows: %d\n", rows);
//ts->printf(cvtest::TS::CONSOLE, "Cols: %d\n", cols);
//ts->printf(cvtest::TS::CONSOLE, "NumChannels: %d\n", num_channels);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
}
}
void CV_MergeTest::run(int)
{
can_merge(1, 1);
can_merge(1, 7);
can_merge(53, 7);
can_merge_submatrixes(1, 1);
can_merge_submatrixes(1, 7);
can_merge_submatrixes(53, 7);
}
////////////////////////////////////////////////////////////////////////////////
// Split
struct CV_SplitTest : public cvtest::BaseTest
{
void can_split(size_t rows, size_t cols);
void can_split_submatrix(size_t rows, size_t cols);
void run(int);
};
void CV_SplitTest::can_split(size_t rows, size_t cols)
{
bool double_ok = gpu::TargetArchs::builtWith(gpu::NATIVE_DOUBLE) &&
gpu::DeviceInfo().supports(gpu::NATIVE_DOUBLE);
size_t depth_end = double_ok ? CV_64F : CV_32F;
for (size_t num_channels = 1; num_channels <= 4; ++num_channels)
for (size_t depth = CV_8U; depth <= depth_end; ++depth)
{
Mat src(rows, cols, CV_MAKETYPE(depth, num_channels), Scalar(1.0, 2.0, 3.0, 4.0));
vector<Mat> dst;
cv::split(src, dst);
gpu::GpuMat dev_src(src);
vector<gpu::GpuMat> dev_dst;
cv::gpu::split(dev_src, dev_dst);
if (dev_dst.size() != dst.size())
{
ts->printf(cvtest::TS::CONSOLE, "Bad output sizes");
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
}
for (size_t i = 0; i < num_channels; ++i)
{
Mat host_dst = dev_dst[i];
double err = norm(dst[i], host_dst, NORM_INF);
if (err > 1e-3)
{
//ts->printf(cvtest::TS::CONSOLE, "\nNorm: %f\n", err);
//ts->printf(cvtest::TS::CONSOLE, "Depth: %d\n", depth);
//ts->printf(cvtest::TS::CONSOLE, "Rows: %d\n", rows);
//ts->printf(cvtest::TS::CONSOLE, "Cols: %d\n", cols);
//ts->printf(cvtest::TS::CONSOLE, "NumChannels: %d\n", num_channels);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
}
}
}
void CV_SplitTest::can_split_submatrix(size_t rows, size_t cols)
{
bool double_ok = gpu::TargetArchs::builtWith(gpu::NATIVE_DOUBLE) &&
gpu::DeviceInfo().supports(gpu::NATIVE_DOUBLE);
size_t depth_end = double_ok ? CV_64F : CV_32F;
for (size_t num_channels = 1; num_channels <= 4; ++num_channels)
for (size_t depth = CV_8U; depth <= depth_end; ++depth)
{
Mat src_data(rows * 2, cols * 2, CV_MAKETYPE(depth, num_channels), Scalar(1.0, 2.0, 3.0, 4.0));
Mat src(src_data(Range(rows / 2, rows / 2 + rows), Range(cols / 2, cols / 2 + cols)));
vector<Mat> dst;
cv::split(src, dst);
gpu::GpuMat dev_src(src);
vector<gpu::GpuMat> dev_dst;
cv::gpu::split(dev_src, dev_dst);
if (dev_dst.size() != dst.size())
{
ts->printf(cvtest::TS::CONSOLE, "Bad output sizes");
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
}
for (size_t i = 0; i < num_channels; ++i)
{
Mat host_dst = dev_dst[i];
double err = norm(dst[i], host_dst, NORM_INF);
if (err > 1e-3)
{
//ts->printf(cvtest::TS::CONSOLE, "\nNorm: %f\n", err);
//ts->printf(cvtest::TS::CONSOLE, "Depth: %d\n", depth);
//ts->printf(cvtest::TS::CONSOLE, "Rows: %d\n", rows);
//ts->printf(cvtest::TS::CONSOLE, "Cols: %d\n", cols);
//ts->printf(cvtest::TS::CONSOLE, "NumChannels: %d\n", num_channels);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
}
}
}
void CV_SplitTest::run(int)
{
can_split(1, 1);
can_split(1, 7);
can_split(7, 53);
can_split_submatrix(1, 1);
can_split_submatrix(1, 7);
can_split_submatrix(7, 53);
}
////////////////////////////////////////////////////////////////////////////////
// Split and merge
struct CV_SplitMergeTest : public cvtest::BaseTest
{
void can_split_merge(size_t rows, size_t cols);
void run(int);
};
void CV_SplitMergeTest::can_split_merge(size_t rows, size_t cols) {
bool double_ok = gpu::TargetArchs::builtWith(gpu::NATIVE_DOUBLE) &&
gpu::DeviceInfo().supports(gpu::NATIVE_DOUBLE);
size_t depth_end = double_ok ? CV_64F : CV_32F;
for (size_t num_channels = 1; num_channels <= 4; ++num_channels)
for (size_t depth = CV_8U; depth <= depth_end; ++depth)
{
Mat orig(rows, cols, CV_MAKETYPE(depth, num_channels), Scalar(1.0, 2.0, 3.0, 4.0));
gpu::GpuMat dev_orig(orig);
vector<gpu::GpuMat> dev_vec;
cv::gpu::split(dev_orig, dev_vec);
gpu::GpuMat dev_final(rows, cols, CV_MAKETYPE(depth, num_channels));
cv::gpu::merge(dev_vec, dev_final);
double err = cv::norm((Mat)dev_orig, (Mat)dev_final, NORM_INF);
if (err > 1e-3)
{
//ts->printf(cvtest::TS::CONSOLE, "\nNorm: %f\n", err);
//ts->printf(cvtest::TS::CONSOLE, "Depth: %d\n", depth);
//ts->printf(cvtest::TS::CONSOLE, "Rows: %d\n", rows);
//ts->printf(cvtest::TS::CONSOLE, "Cols: %d\n", cols);
//ts->printf(cvtest::TS::CONSOLE, "NumChannels: %d\n", num_channels);
ts->set_failed_test_info(cvtest::TS::FAIL_INVALID_OUTPUT);
return;
}
}
}
void CV_SplitMergeTest::run(int)
{
can_split_merge(1, 1);
can_split_merge(1, 7);
can_split_merge(7, 53);
}
TEST(merge, accuracy) { CV_MergeTest test; test.safe_run(); }
TEST(split, accuracy) { CV_SplitTest test; test.safe_run(); }
TEST(split, merge_consistency) { CV_SplitMergeTest test; test.safe_run(); }

131
modules/gpu/test/test_stereo_bm.cpp
View File

@ -1,131 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
using namespace cv;
using namespace cv::gpu;
struct CV_GpuStereoBMTest : public cvtest::BaseTest
{
void run_stress()
{
RNG rng;
for(int i = 0; i < 10; ++i)
{
int winSize = cvRound(rng.uniform(2, 11)) * 2 + 1;
for(int j = 0; j < 10; ++j)
{
int ndisp = cvRound(rng.uniform(5, 32)) * 8;
for(int s = 0; s < 10; ++s)
{
int w = cvRound(rng.uniform(1024, 2048));
int h = cvRound(rng.uniform(768, 1152));
for(int p = 0; p < 2; ++p)
{
//int winSize = winsz[i];
//int disp = disps[j];
Size imgSize(w, h);//res[s];
int preset = p;
printf("Preset = %d, nidsp = %d, winsz = %d, width = %d, height = %d\n", p, ndisp, winSize, imgSize.width, imgSize.height);
GpuMat l(imgSize, CV_8U);
GpuMat r(imgSize, CV_8U);
GpuMat disparity;
StereoBM_GPU bm(preset, ndisp, winSize);
bm(l, r, disparity);
}
}
}
}
}
void run(int )
{
/*run_stress();
return;*/
cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-L.png", 0);
cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-R.png", 0);
cv::Mat img_reference = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", 0);
if (img_l.empty() || img_r.empty() || img_reference.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
cv::gpu::GpuMat disp;
cv::gpu::StereoBM_GPU bm(0, 128, 19);
bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
disp.convertTo(disp, img_reference.type());
double norm = cv::norm((Mat)disp, img_reference, cv::NORM_INF);
//cv::imwrite(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", disp);
/*cv::imshow("disp", disp);
cv::imshow("img_reference", img_reference);
cv::Mat diff = (cv::Mat)disp - (cv::Mat)img_reference;
cv::imshow("diff", diff);
cv::waitKey();*/
if (norm >= 100)
{
ts->printf(cvtest::TS::LOG, "\nStereoBM norm = %f\n", norm);
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
ts->set_failed_test_info(cvtest::TS::OK);
}
};
TEST(StereoBM, regression) { CV_GpuStereoBMTest test; test.safe_run(); }

86
modules/gpu/test/test_stereo_bm_async.cpp
View File

@ -1,86 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
using namespace cv;
using namespace std;
struct CV_AsyncStereoBMTest : public cvtest::BaseTest
{
void run( int /* start_from */)
{
cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-L.png", 0);
cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-R.png", 0);
cv::Mat img_reference = cv::imread(std::string(ts->get_data_path()) + "stereobm/aloe-disp.png", 0);
if (img_l.empty() || img_r.empty() || img_reference.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
cv::gpu::GpuMat disp;
cv::gpu::StereoBM_GPU bm(0, 128, 19);
cv::gpu::Stream stream;
for (size_t i = 0; i < 50; i++)
{
bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp, stream);
}
stream.waitForCompletion();
disp.convertTo(disp, img_reference.type());
double norm = cv::norm((Mat)disp, img_reference, cv::NORM_INF);
if (norm >= 100)
{
ts->printf(cvtest::TS::LOG, "\nStereoBM norm = %f\n", norm);
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
ts->set_failed_test_info(cvtest::TS::OK);
}
};
TEST(StereoBM, async) { CV_AsyncStereoBMTest test; test.safe_run(); }

82
modules/gpu/test/test_stereo_bp.cpp
View File

@ -1,82 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
struct CV_GpuStereoBPTest : public cvtest::BaseTest
{
void run(int)
{
cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-L.png");
cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-R.png");
cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-disp.png", 0);
if (img_l.empty() || img_r.empty() || img_template.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
{cv::Mat temp; cv::cvtColor(img_l, temp, CV_BGR2BGRA); cv::swap(temp, img_l);}
{cv::Mat temp; cv::cvtColor(img_r, temp, CV_BGR2BGRA); cv::swap(temp, img_r);}
cv::gpu::StereoBeliefPropagation bpm(64, 8, 2, 25, 0.1f, 15, 1, CV_16S);
cv::gpu::GpuMat disp;
bpm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
//cv::imwrite(std::string(ts->get_data_path()) + "stereobp/aloe-disp.png", disp);
disp.convertTo(disp, img_template.type());
double norm = cv::norm((cv::Mat)disp, img_template, cv::NORM_INF);
if (norm >= 0.5)
{
ts->printf(cvtest::TS::LOG, "\nStereoBP norm = %f\n", norm);
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
ts->set_failed_test_info(cvtest::TS::OK);
}
};
TEST(StereoBP, regression) { CV_GpuStereoBPTest test; test.safe_run(); }

89
modules/gpu/test/test_stereo_csbp.cpp
View File

@ -1,89 +0,0 @@
/*M///////////////////////////////////////////////////////////////////////////////////////
//
// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
//
// By downloading, copying, installing or using the software you agree to this license.
// If you do not agree to this license, do not download, install,
// copy or use the software.
//
//
// Intel License Agreement
// For Open Source Computer Vision Library
//
// Copyright (C) 2000, Intel Corporation, all rights reserved.
// Third party copyrights are property of their respective owners.
//
// Redistribution and use in source and binary forms, with or without modification,
// are permitted provided that the following conditions are met:
//
// * Redistribution's of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
//
// * Redistribution's in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// * The name of Intel Corporation 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"
struct CV_GpuStereoCSBPTest : public cvtest::BaseTest
{
void run(int )
{
cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-L.png");
cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-R.png");
cv::Mat img_template;
if (cv::gpu::TargetArchs::builtWith(cv::gpu::FEATURE_SET_COMPUTE_20) &&
cv::gpu::DeviceInfo().supports(cv::gpu::FEATURE_SET_COMPUTE_20))
img_template = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-disp.png", CV_LOAD_IMAGE_GRAYSCALE);
else
img_template = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-disp_CC1X.png", CV_LOAD_IMAGE_GRAYSCALE);
if (img_l.empty() || img_r.empty() || img_template.empty())
{
ts->set_failed_test_info(cvtest::TS::FAIL_MISSING_TEST_DATA);
return;
}
{cv::Mat temp; cv::cvtColor(img_l, temp, CV_BGR2BGRA); cv::swap(temp, img_l);}
{cv::Mat temp; cv::cvtColor(img_r, temp, CV_BGR2BGRA); cv::swap(temp, img_r);}
cv::gpu::GpuMat disp;
cv::gpu::StereoConstantSpaceBP bpm(128, 16, 4, 4);
bpm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
//cv::imwrite(std::string(ts->get_data_path()) + "csstereobp/aloe-disp_CC1X.png", cv::Mat(disp));
disp.convertTo(disp, img_template.type());
double norm = cv::norm((cv::Mat)disp, img_template, cv::NORM_INF);
if (norm >= 1.5)
{
ts->printf(cvtest::TS::LOG, "\nConstantSpaceStereoBP norm = %f\n", norm);
ts->set_failed_test_info(cvtest::TS::FAIL_GENERIC);
return;
}
ts->set_failed_test_info(cvtest::TS::OK);
}
};
TEST(StereoConstantSpaceBP, regression) { CV_GpuStereoCSBPTest test; test.safe_run(); }