Commit 059cef57 authored by Vladislav Vinogradov's avatar Vladislav Vinogradov

fixed gpu::filter2D border interpolation for CV_32FC1 type

added additional tests for gpu filters
fixed gpu features2D tests
parent c1a6cb62
......@@ -1661,7 +1661,7 @@ public:
};
//! Constructor
explicit ORB_GPU(int nFeatures = 500, float scaleFactor = 1.2f, int nLevels = 3, int edgeThreshold = 31,
explicit ORB_GPU(int nFeatures = 500, float scaleFactor = 1.2f, int nLevels = 8, int edgeThreshold = 31,
int firstLevel = 0, int WTA_K = 2, int scoreType = 0, int patchSize = 31);
//! Compute the ORB features on an image
......
......@@ -911,27 +911,28 @@ namespace cv { namespace gpu { namespace device
__constant__ float c_filter2DKernel[FILTER2D_MAX_KERNEL_SIZE * FILTER2D_MAX_KERNEL_SIZE];
texture<float, cudaTextureType2D, cudaReadModeElementType> filter2DTex(0, cudaFilterModePoint, cudaAddressModeBorder);
texture<float, cudaTextureType2D, cudaReadModeElementType> filter2DTex(0, cudaFilterModePoint, cudaAddressModeClamp);
__global__ void filter2D(int ofsX, int ofsY, DevMem2Df dst, const int kWidth, const int kHeight, const int anchorX, const int anchorY)
__global__ void filter2D(int ofsX, int ofsY, PtrStepf dst, const int kWidth, const int kHeight, const int anchorX, const int anchorY, const BrdReflect101<float> brd)
{
const int x = blockIdx.x * blockDim.x + threadIdx.x;
const int y = blockIdx.y * blockDim.y + threadIdx.y;
if (x >= dst.cols || y >= dst.rows)
if (x > brd.last_col || y > brd.last_row)
return;
float res = 0;
const int baseX = ofsX + x - anchorX;
const int baseY = ofsY + y - anchorY;
int kInd = 0;
for (int i = 0; i < kHeight; ++i)
{
for (int j = 0; j < kWidth; ++j)
res += tex2D(filter2DTex, baseX + j, baseY + i) * c_filter2DKernel[kInd++];
{
const int srcX = ofsX + brd.idx_col(x - anchorX + j);
const int srcY = ofsY + brd.idx_row(y - anchorY + i);
res += tex2D(filter2DTex, srcX, srcY) * c_filter2DKernel[kInd++];
}
}
dst.ptr(y)[x] = res;
......@@ -946,7 +947,9 @@ namespace cv { namespace gpu { namespace device
bindTexture(&filter2DTex, src);
filter2D<<<grid, block, 0, stream>>>(ofsX, ofsY, dst, kWidth, kHeight, anchorX, anchorY);
BrdReflect101<float> brd(dst.rows, dst.cols);
filter2D<<<grid, block, 0, stream>>>(ofsX, ofsY, dst, kWidth, kHeight, anchorX, anchorY, brd);
cudaSafeCall(cudaGetLastError());
if (stream == 0)
......
......@@ -238,7 +238,7 @@ namespace
cv::gpu::SURF_GPU::SURF_GPU()
{
hessianThreshold = 100;
extended = 1;
extended = true;
nOctaves = 4;
nOctaveLayers = 2;
keypointsRatio = 0.01f;
......
This diff is collapsed.
......@@ -41,9 +41,11 @@
#include "precomp.hpp"
#ifdef HAVE_CUDA
namespace {
PARAM_TEST_CASE(CopyMakeBorder, cv::gpu::DeviceInfo, cv::Size, MatType, int, Border, UseRoi)
IMPLEMENT_PARAM_CLASS(Border, int)
PARAM_TEST_CASE(CopyMakeBorder, cv::gpu::DeviceInfo, cv::Size, MatType, Border, BorderType, UseRoi)
{
cv::gpu::DeviceInfo devInfo;
cv::Size size;
......@@ -82,9 +84,17 @@ TEST_P(CopyMakeBorder, Accuracy)
INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CopyMakeBorder, testing::Combine(
ALL_DEVICES,
DIFFERENT_SIZES,
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)),
testing::Values(1, 10, 50),
testing::Values(Border(cv::BORDER_REFLECT101), Border(cv::BORDER_REPLICATE), Border(cv::BORDER_CONSTANT), Border(cv::BORDER_REFLECT), Border(cv::BORDER_WRAP)),
testing::Values(MatType(CV_8UC1),
MatType(CV_8UC3),
MatType(CV_8UC4),
MatType(CV_16UC1),
MatType(CV_16UC3),
MatType(CV_16UC4),
MatType(CV_32FC1),
MatType(CV_32FC3),
MatType(CV_32FC4)),
testing::Values(Border(1), Border(10), Border(50)),
ALL_BORDER_TYPES,
WHOLE_SUBMAT));
#endif // HAVE_CUDA
} // namespace
This diff is collapsed.
This diff is collapsed.
This diff is collapsed.
......@@ -2198,7 +2198,7 @@ INSTANTIATE_TEST_CASE_P(ImgProc, EqualizeHist, ALL_DEVICES);
///////////////////////////////////////////////////////////////////////////////////////////////////////
// cornerHarris
PARAM_TEST_CASE(CornerHarris, cv::gpu::DeviceInfo, MatType, Border, int, int)
PARAM_TEST_CASE(CornerHarris, cv::gpu::DeviceInfo, MatType, BorderType, int, int)
{
cv::gpu::DeviceInfo devInfo;
int type;
......@@ -2257,7 +2257,7 @@ INSTANTIATE_TEST_CASE_P(ImgProc, CornerHarris, Combine(
///////////////////////////////////////////////////////////////////////////////////////////////////////
// cornerMinEigen
PARAM_TEST_CASE(CornerMinEigen, cv::gpu::DeviceInfo, MatType, Border, int, int)
PARAM_TEST_CASE(CornerMinEigen, cv::gpu::DeviceInfo, MatType, BorderType, int, int)
{
cv::gpu::DeviceInfo devInfo;
int type;
......@@ -2572,6 +2572,8 @@ INSTANTIATE_TEST_CASE_P(ImgProc, MeanShiftSegmentation, Combine(
////////////////////////////////////////////////////////////////////////////////
// matchTemplate
CV_ENUM(TemplateMethod, cv::TM_SQDIFF, cv::TM_SQDIFF_NORMED, cv::TM_CCORR, cv::TM_CCORR_NORMED, cv::TM_CCOEFF, cv::TM_CCOEFF_NORMED)
PARAM_TEST_CASE(MatchTemplate8U, cv::gpu::DeviceInfo, int, TemplateMethod)
{
cv::gpu::DeviceInfo devInfo;
......@@ -2776,6 +2778,8 @@ INSTANTIATE_TEST_CASE_P(ImgProc, MatchTemplate_CCOEF_NORMED, Combine(
////////////////////////////////////////////////////////////////////////////
// MulSpectrums
CV_FLAGS(DftFlags, cv::DFT_INVERSE, cv::DFT_SCALE, cv::DFT_ROWS, cv::DFT_COMPLEX_OUTPUT, cv::DFT_REAL_OUTPUT)
PARAM_TEST_CASE(MulSpectrums, cv::gpu::DeviceInfo, DftFlags)
{
cv::gpu::DeviceInfo devInfo;
......
......@@ -109,7 +109,7 @@ namespace
///////////////////////////////////////////////////////////////////
// Test
PARAM_TEST_CASE(Remap, cv::gpu::DeviceInfo, cv::Size, MatType, Interpolation, Border, UseRoi)
PARAM_TEST_CASE(Remap, cv::gpu::DeviceInfo, cv::Size, MatType, Interpolation, BorderType, UseRoi)
{
cv::gpu::DeviceInfo devInfo;
cv::Size size;
......@@ -171,7 +171,7 @@ INSTANTIATE_TEST_CASE_P(GPU_ImgProc, Remap, testing::Combine(
DIFFERENT_SIZES,
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)),
testing::Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)),
testing::Values(Border(cv::BORDER_REFLECT101), Border(cv::BORDER_REPLICATE), Border(cv::BORDER_CONSTANT), Border(cv::BORDER_REFLECT), Border(cv::BORDER_WRAP)),
testing::Values(BorderType(cv::BORDER_REFLECT101), BorderType(cv::BORDER_REPLICATE), BorderType(cv::BORDER_CONSTANT), BorderType(cv::BORDER_REFLECT), BorderType(cv::BORDER_WRAP)),
WHOLE_SUBMAT));
#endif // HAVE_CUDA
......@@ -43,6 +43,9 @@
#ifdef HAVE_CUDA
CV_ENUM(ThreshOp, cv::THRESH_BINARY, cv::THRESH_BINARY_INV, cv::THRESH_TRUNC, cv::THRESH_TOZERO, cv::THRESH_TOZERO_INV)
#define ALL_THRESH_OPS testing::Values(ThreshOp(cv::THRESH_BINARY), ThreshOp(cv::THRESH_BINARY_INV), ThreshOp(cv::THRESH_TRUNC), ThreshOp(cv::THRESH_TOZERO), ThreshOp(cv::THRESH_TOZERO_INV))
PARAM_TEST_CASE(Threshold, cv::gpu::DeviceInfo, cv::Size, MatType, ThreshOp, UseRoi)
{
cv::gpu::DeviceInfo devInfo;
......
......@@ -175,7 +175,7 @@ namespace
///////////////////////////////////////////////////////////////////
// Test
PARAM_TEST_CASE(WarpAffine, cv::gpu::DeviceInfo, cv::Size, MatType, Inverse, Interpolation, Border, UseRoi)
PARAM_TEST_CASE(WarpAffine, cv::gpu::DeviceInfo, cv::Size, MatType, Inverse, Interpolation, BorderType, UseRoi)
{
cv::gpu::DeviceInfo devInfo;
cv::Size size;
......@@ -225,7 +225,7 @@ INSTANTIATE_TEST_CASE_P(GPU_ImgProc, WarpAffine, testing::Combine(
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)),
DIRECT_INVERSE,
testing::Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)),
testing::Values(Border(cv::BORDER_REFLECT101), Border(cv::BORDER_REPLICATE), Border(cv::BORDER_REFLECT), Border(cv::BORDER_WRAP)),
testing::Values(BorderType(cv::BORDER_REFLECT101), BorderType(cv::BORDER_REPLICATE), BorderType(cv::BORDER_REFLECT), BorderType(cv::BORDER_WRAP)),
WHOLE_SUBMAT));
///////////////////////////////////////////////////////////////////
......
......@@ -175,7 +175,7 @@ namespace
///////////////////////////////////////////////////////////////////
// Test
PARAM_TEST_CASE(WarpPerspective, cv::gpu::DeviceInfo, cv::Size, MatType, Inverse, Interpolation, Border, UseRoi)
PARAM_TEST_CASE(WarpPerspective, cv::gpu::DeviceInfo, cv::Size, MatType, Inverse, Interpolation, BorderType, UseRoi)
{
cv::gpu::DeviceInfo devInfo;
cv::Size size;
......@@ -225,7 +225,7 @@ INSTANTIATE_TEST_CASE_P(GPU_ImgProc, WarpPerspective, testing::Combine(
testing::Values(MatType(CV_8UC1), MatType(CV_8UC3), MatType(CV_8UC4), MatType(CV_16UC1), MatType(CV_16UC3), MatType(CV_16UC4), MatType(CV_32FC1), MatType(CV_32FC3), MatType(CV_32FC4)),
DIRECT_INVERSE,
testing::Values(Interpolation(cv::INTER_NEAREST), Interpolation(cv::INTER_LINEAR), Interpolation(cv::INTER_CUBIC)),
testing::Values(Border(cv::BORDER_REFLECT101), Border(cv::BORDER_REPLICATE), Border(cv::BORDER_REFLECT), Border(cv::BORDER_WRAP)),
testing::Values(BorderType(cv::BORDER_REFLECT101), BorderType(cv::BORDER_REPLICATE), BorderType(cv::BORDER_REFLECT), BorderType(cv::BORDER_WRAP)),
WHOLE_SUBMAT));
///////////////////////////////////////////////////////////////////
......
......@@ -47,6 +47,9 @@ using namespace cv::gpu;
using namespace cvtest;
using namespace testing;
//////////////////////////////////////////////////////////////////////
// random generators
int randomInt(int minVal, int maxVal)
{
RNG& rng = TS::ptr()->get_rng();
......@@ -74,6 +77,9 @@ Mat randomMat(Size size, int type, double minVal, double maxVal)
return randomMat(TS::ptr()->get_rng(), size, type, minVal, maxVal, false);
}
//////////////////////////////////////////////////////////////////////
// GpuMat create
cv::gpu::GpuMat createMat(cv::Size size, int type, bool useRoi)
{
Size size0 = size;
......@@ -99,6 +105,30 @@ GpuMat loadMat(const Mat& m, bool useRoi)
return d_m;
}
//////////////////////////////////////////////////////////////////////
// Image load
Mat readImage(const string& fileName, int flags)
{
return imread(string(cvtest::TS::ptr()->get_data_path()) + fileName, flags);
}
Mat readImageType(const string& fname, int type)
{
Mat src = readImage(fname, CV_MAT_CN(type) == 1 ? IMREAD_GRAYSCALE : IMREAD_COLOR);
if (CV_MAT_CN(type) == 4)
{
Mat temp;
cvtColor(src, temp, cv::COLOR_BGR2BGRA);
swap(src, temp);
}
src.convertTo(src, CV_MAT_DEPTH(type));
return src;
}
//////////////////////////////////////////////////////////////////////
// Gpu devices
bool supportFeature(const DeviceInfo& info, FeatureSet feature)
{
return TargetArchs::builtWith(feature) && info.supports(feature);
......@@ -150,86 +180,146 @@ vector<DeviceInfo> devices(FeatureSet feature)
return devs_filtered;
}
vector<MatType> types(int depth_start, int depth_end, int cn_start, int cn_end)
{
vector<MatType> v;
v.reserve((depth_end - depth_start + 1) * (cn_end - cn_start + 1));
//////////////////////////////////////////////////////////////////////
// Additional assertion
for (int depth = depth_start; depth <= depth_end; ++depth)
{
for (int cn = cn_start; cn <= cn_end; ++cn)
Mat getMat(InputArray arr)
{
if (arr.kind() == _InputArray::GPU_MAT)
{
v.push_back(CV_MAKETYPE(depth, cn));
}
Mat m;
arr.getGpuMat().download(m);
return m;
}
return v;
return arr.getMat();
}
const vector<MatType>& all_types()
double checkNorm(InputArray m1, const InputArray m2)
{
static vector<MatType> v = types(CV_8U, CV_64F, 1, 4);
return v;
return norm(getMat(m1), getMat(m2), NORM_INF);
}
Mat readImage(const string& fileName, int flags)
void minMaxLocGold(const Mat& src, double* minVal_, double* maxVal_, Point* minLoc_, Point* maxLoc_, const Mat& mask)
{
return imread(string(cvtest::TS::ptr()->get_data_path()) + fileName, flags);
}
if (src.depth() != CV_8S)
{
minMaxLoc(src, minVal_, maxVal_, minLoc_, maxLoc_, mask);
return;
}
Mat readImageType(const string& fname, int type)
{
Mat src = readImage(fname, CV_MAT_CN(type) == 1 ? IMREAD_GRAYSCALE : IMREAD_COLOR);
if (CV_MAT_CN(type) == 4)
// OpenCV's minMaxLoc doesn't support CV_8S type
double minVal = numeric_limits<double>::max();
Point minLoc(-1, -1);
double maxVal = -numeric_limits<double>::max();
Point maxLoc(-1, -1);
for (int y = 0; y < src.rows; ++y)
{
Mat temp;
cvtColor(src, temp, cv::COLOR_BGR2BGRA);
swap(src, temp);
const schar* src_row = src.ptr<signed char>(y);
const uchar* mask_row = mask.empty() ? 0 : mask.ptr<unsigned char>(y);
for (int x = 0; x < src.cols; ++x)
{
if (!mask_row || mask_row[x])
{
schar val = src_row[x];
if (val < minVal)
{
minVal = val;
minLoc = cv::Point(x, y);
}
src.convertTo(src, CV_MAT_DEPTH(type));
return src;
if (val > maxVal)
{
maxVal = val;
maxLoc = cv::Point(x, y);
}
}
}
}
if (minVal_) *minVal_ = minVal;
if (maxVal_) *maxVal_ = maxVal;
if (minLoc_) *minLoc_ = minLoc;
if (maxLoc_) *maxLoc_ = maxLoc;
}
namespace
{
Mat getMat(InputArray arr)
template <typename T, typename OutT> string printMatValImpl(const Mat& m, Point p)
{
if (arr.kind() == _InputArray::GPU_MAT)
const int cn = m.channels();
ostringstream ostr;
ostr << "(";
p.x /= cn;
ostr << static_cast<OutT>(m.at<T>(p.y, p.x * cn));
for (int c = 1; c < m.channels(); ++c)
{
Mat m;
arr.getGpuMat().download(m);
return m;
ostr << ", " << static_cast<OutT>(m.at<T>(p.y, p.x * cn + c));
}
ostr << ")";
return arr.getMat();
return ostr.str();
}
string printMatVal(const Mat& m, Point p)
{
typedef string (*func_t)(const Mat& m, Point p);
static const func_t funcs[] =
{
printMatValImpl<uchar, int>, printMatValImpl<schar, int>, printMatValImpl<ushort, int>, printMatValImpl<short, int>,
printMatValImpl<int, int>, printMatValImpl<float, float>, printMatValImpl<double, double>
};
return funcs[m.depth()](m, p);
}
}
void showDiff(InputArray gold_, InputArray actual_, double eps)
testing::AssertionResult assertMatNear(const char* expr1, const char* expr2, const char* eps_expr, cv::InputArray m1_, cv::InputArray m2_, double eps)
{
Mat gold = getMat(gold_);
Mat actual = getMat(actual_);
Mat m1 = getMat(m1_);
Mat m2 = getMat(m2_);
Mat diff;
absdiff(gold, actual, diff);
threshold(diff, diff, eps, 255.0, cv::THRESH_BINARY);
if (m1.size() != m2.size())
{
return AssertionFailure() << "Matrices \"" << expr1 << "\" and \"" << expr2 << "\" have different sizes : \""
<< expr1 << "\" [" << PrintToString(m1.size()) << "] vs \""
<< expr2 << "\" [" << PrintToString(m2.size()) << "]";
}
namedWindow("gold", WINDOW_NORMAL);
namedWindow("actual", WINDOW_NORMAL);
namedWindow("diff", WINDOW_NORMAL);
if (m1.type() != m2.type())
{
return AssertionFailure() << "Matrices \"" << expr1 << "\" and \"" << expr2 << "\" have different types : \""
<< expr1 << "\" [" << PrintToString(MatType(m1.type())) << "] vs \""
<< expr2 << "\" [" << PrintToString(MatType(m2.type())) << "]";
}
imshow("gold", gold);
imshow("actual", actual);
imshow("diff", diff);
Mat diff;
absdiff(m1.reshape(1), m2.reshape(1), diff);
waitKey();
}
double maxVal = 0.0;
Point maxLoc;
minMaxLocGold(diff, 0, &maxVal, 0, &maxLoc);
double checkNorm(InputArray m1, const InputArray m2)
{
return norm(getMat(m1), getMat(m2), NORM_INF);
if (maxVal > eps)
{
return AssertionFailure() << "The max difference between matrices \"" << expr1 << "\" and \"" << expr2
<< "\" is " << maxVal << " at (" << maxLoc.y << ", " << maxLoc.x / m1.channels() << ")"
<< ", which exceeds \"" << eps_expr << "\", where \""
<< expr1 << "\" at (" << maxLoc.y << ", " << maxLoc.x / m1.channels() << ") evaluates to " << printMatVal(m1, maxLoc) << ", \""
<< expr2 << "\" at (" << maxLoc.y << ", " << maxLoc.x / m1.channels() << ") evaluates to " << printMatVal(m2, maxLoc) << ", \""
<< eps_expr << "\" evaluates to " << eps;
}
return AssertionSuccess();
}
double checkSimilarity(InputArray m1, InputArray m2)
......@@ -239,6 +329,45 @@ double checkSimilarity(InputArray m1, InputArray m2)
return std::abs(diff.at<float>(0, 0) - 1.f);
}
//////////////////////////////////////////////////////////////////////
// Helper structs for value-parameterized tests
vector<MatDepth> depths(int depth_start, int depth_end)
{
vector<MatDepth> v;
v.reserve((depth_end - depth_start + 1));
for (int depth = depth_start; depth <= depth_end; ++depth)
v.push_back(depth);
return v;
}
vector<MatType> types(int depth_start, int depth_end, int cn_start, int cn_end)
{
vector<MatType> 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 vector<MatType>& all_types()
{
static vector<MatType> v = types(CV_8U, CV_64F, 1, 4);
return v;
}
void cv::gpu::PrintTo(const DeviceInfo& info, ostream* os)
{
(*os) << info.name();
......@@ -259,3 +388,23 @@ void PrintTo(const Inverse& inverse, std::ostream* os)
else
(*os) << "direct";
}
void showDiff(InputArray gold_, InputArray actual_, double eps)
{
Mat gold = getMat(gold_);
Mat actual = getMat(actual_);
Mat diff;
absdiff(gold, actual, diff);
threshold(diff, diff, eps, 255.0, cv::THRESH_BINARY);
namedWindow("gold", WINDOW_NORMAL);
namedWindow("actual", WINDOW_NORMAL);
namedWindow("diff", WINDOW_NORMAL);
imshow("gold", gold);
imshow("actual", actual);
imshow("diff", diff);
waitKey();
}
This diff is collapsed.
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment