cosmetic changes in gpu module, decreased matchTemplate test running time

pull/13383/head
Alexey Spizhevoy 14 years ago
parent d091ae5746
commit 4dfbf99dd5
  1. 4
      modules/gpu/include/opencv2/gpu/gpu.hpp
  2. 20
      modules/gpu/src/cuda/match_template.cu
  3. 10
      modules/gpu/src/hog.cpp
  4. 4
      samples/gpu/hog.cpp
  5. 48
      tests/gpu/src/match_template.cpp

@ -1182,8 +1182,8 @@ namespace cv
void setSVMDetector(const vector<float>& detector);
static vector<float> getDefaultPeopleDetector();
static vector<float> getPeopleDetector_48x96();
static vector<float> getPeopleDetector_64x128();
static vector<float> getPeopleDetector48x96();
static vector<float> getPeopleDetector64x128();
void detect(const GpuMat& img, vector<Point>& found_locations,
double hit_threshold=0, Size win_stride=Size(),

@ -560,7 +560,7 @@ __global__ void matchTemplatePreparedKernel_CCOFF_NORMED_8U(
(image_sqsum.ptr(y + h)[x + w] - image_sqsum.ptr(y)[x + w]) -
(image_sqsum.ptr(y + h)[x] - image_sqsum.ptr(y)[x]));
result.ptr(y)[x] = (ccorr - image_sum_ * templ_sum_scale) *
rsqrtf(templ_sqsum_scale * max(1.f, image_sqsum_ - weight * image_sum_ * image_sum_));
rsqrtf(templ_sqsum_scale * max(1e-3f, image_sqsum_ - weight * image_sum_ * image_sum_));
}
}
@ -610,8 +610,8 @@ __global__ void matchTemplatePreparedKernel_CCOFF_NORMED_8UC2(
(image_sqsum_g.ptr(y + h)[x + w] - image_sqsum_g.ptr(y)[x + w]) -
(image_sqsum_g.ptr(y + h)[x] - image_sqsum_g.ptr(y)[x]));
float ccorr = result.ptr(y)[x];
float rdenom = rsqrtf(templ_sqsum_scale * max(1.f, image_sqsum_r_ - weight * image_sum_r_ * image_sum_r_
+ image_sqsum_g_ - weight * image_sum_g_ * image_sum_g_));
float rdenom = rsqrtf(templ_sqsum_scale * max(1e-3f, image_sqsum_r_ - weight * image_sum_r_ * image_sum_r_
+ image_sqsum_g_ - weight * image_sum_g_ * image_sum_g_));
result.ptr(y)[x] = (ccorr - image_sum_r_ * templ_sum_scale_r
- image_sum_g_ * templ_sum_scale_g) * rdenom;
}
@ -678,9 +678,9 @@ __global__ void matchTemplatePreparedKernel_CCOFF_NORMED_8UC3(
(image_sqsum_b.ptr(y + h)[x + w] - image_sqsum_b.ptr(y)[x + w]) -
(image_sqsum_b.ptr(y + h)[x] - image_sqsum_b.ptr(y)[x]));
float ccorr = result.ptr(y)[x];
float rdenom = rsqrtf(templ_sqsum_scale * max(1.f, image_sqsum_r_ - weight * image_sum_r_ * image_sum_r_
+ image_sqsum_g_ - weight * image_sum_g_ * image_sum_g_
+ image_sqsum_b_ - weight * image_sum_b_ * image_sum_b_));
float rdenom = rsqrtf(templ_sqsum_scale * max(1e-3f, image_sqsum_r_ - weight * image_sum_r_ * image_sum_r_
+ image_sqsum_g_ - weight * image_sum_g_ * image_sum_g_
+ image_sqsum_b_ - weight * image_sum_b_ * image_sum_b_));
result.ptr(y)[x] = (ccorr - image_sum_r_ * templ_sum_scale_r
- image_sum_g_ * templ_sum_scale_g
- image_sum_b_ * templ_sum_scale_b) * rdenom;
@ -760,10 +760,10 @@ __global__ void matchTemplatePreparedKernel_CCOFF_NORMED_8UC4(
(image_sqsum_a.ptr(y + h)[x + w] - image_sqsum_a.ptr(y)[x + w]) -
(image_sqsum_a.ptr(y + h)[x] - image_sqsum_a.ptr(y)[x]));
float ccorr = result.ptr(y)[x];
float rdenom = rsqrtf(templ_sqsum_scale * max(1.f, image_sqsum_r_ - weight * image_sum_r_ * image_sum_r_
+ image_sqsum_g_ - weight * image_sum_g_ * image_sum_g_
+ image_sqsum_b_ - weight * image_sum_b_ * image_sum_b_
+ image_sqsum_a_ - weight * image_sum_a_ * image_sum_a_));
float rdenom = rsqrtf(templ_sqsum_scale * max(1e-3f, image_sqsum_r_ - weight * image_sum_r_ * image_sum_r_
+ image_sqsum_g_ - weight * image_sum_g_ * image_sum_g_
+ image_sqsum_b_ - weight * image_sum_b_ * image_sum_b_
+ image_sqsum_a_ - weight * image_sum_a_ * image_sum_a_));
result.ptr(y)[x] = (ccorr - image_sum_r_ * templ_sum_scale_r
- image_sum_g_ * templ_sum_scale_g
- image_sum_b_ * templ_sum_scale_b

@ -55,8 +55,8 @@ void cv::gpu::HOGDescriptor::detectMultiScale(const GpuMat&, vector<Rect>&, doub
void cv::gpu::HOGDescriptor::computeBlockHistograms(const GpuMat&) { throw_nogpu(); }
void cv::gpu::HOGDescriptor::getDescriptors(const GpuMat&, Size, GpuMat&, int) { throw_nogpu(); }
std::vector<float> cv::gpu::HOGDescriptor::getDefaultPeopleDetector() { throw_nogpu(); return std::vector<float>(); }
std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector_48x96() { throw_nogpu(); return std::vector<float>(); }
std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector_64x128() { throw_nogpu(); return std::vector<float>(); }
std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector48x96() { throw_nogpu(); return std::vector<float>(); }
std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector64x128() { throw_nogpu(); return std::vector<float>(); }
#else
@ -352,11 +352,11 @@ cv::Size cv::gpu::HOGDescriptor::numPartsWithin(cv::Size size, cv::Size part_siz
std::vector<float> cv::gpu::HOGDescriptor::getDefaultPeopleDetector()
{
return getPeopleDetector_64x128();
return getPeopleDetector64x128();
}
std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector_48x96()
std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector48x96()
{
static const float detector[] = {
0.294350f, -0.098796f, -0.129522f, 0.078753f, 0.387527f, 0.261529f,
@ -696,7 +696,7 @@ std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector_48x96()
std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector_64x128()
std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector64x128()
{
static const float detector[] = {
0.05359386f, -0.14721455f, -0.05532170f, 0.05077307f,

@ -239,9 +239,9 @@ void App::run()
// Create HOG descriptors and detectors here
vector<float> detector;
if (win_size == Size(64, 128))
detector = cv::gpu::HOGDescriptor::getPeopleDetector_64x128();
detector = cv::gpu::HOGDescriptor::getPeopleDetector64x128();
else
detector = cv::gpu::HOGDescriptor::getPeopleDetector_48x96();
detector = cv::gpu::HOGDescriptor::getPeopleDetector48x96();
cv::gpu::HOGDescriptor gpu_hog(win_size, Size(16, 16), Size(8, 8), Size(8, 8), 9,
cv::gpu::HOGDescriptor::DEFAULT_WIN_SIGMA, 0.2, gamma_corr,

@ -70,17 +70,17 @@ struct CV_GpuMatchTemplateTest: CvTest
int n, m, h, w;
F(clock_t t;)
RNG rng(*ts->get_rng());
for (int cn = 1; cn <= 4; ++cn)
{
F(ts->printf(CvTS::CONSOLE, "cn: %d\n", cn);)
for (int i = 0; i <= 0; ++i)
{
n = 1 + rand() % 1000;
m = 1 + rand() % 1000;
do h = 1 + rand() % 100; while (h > n);
do w = 1 + rand() % 100; while (w > m);
//cout << "w: " << w << " h: " << h << endl;
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);
@ -91,7 +91,7 @@ struct CV_GpuMatchTemplateTest: CvTest
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)) return;
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);
@ -102,7 +102,7 @@ struct CV_GpuMatchTemplateTest: CvTest
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)) return;
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);
@ -113,7 +113,7 @@ struct CV_GpuMatchTemplateTest: CvTest
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)) return;
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);
@ -124,7 +124,7 @@ struct CV_GpuMatchTemplateTest: CvTest
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)) return;
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);
@ -135,7 +135,7 @@ struct CV_GpuMatchTemplateTest: CvTest
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)) return;
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);
@ -146,7 +146,7 @@ struct CV_GpuMatchTemplateTest: CvTest
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)) return;
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);
@ -157,7 +157,7 @@ struct CV_GpuMatchTemplateTest: CvTest
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)) return;
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);
@ -168,7 +168,7 @@ struct CV_GpuMatchTemplateTest: CvTest
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)) return;
if (!check(dst_gold, Mat(dst), 0.25f * h * w * 1e-5f, "CCORR 32F")) return;
}
}
}
@ -190,19 +190,33 @@ struct CV_GpuMatchTemplateTest: CvTest
rng.fill(a, RNG::UNIFORM, Scalar::all(0.001f), Scalar::all(1.f));
}
bool check(const Mat& a, const Mat& b, float max_err)
bool check(const Mat& a, const Mat& b, float max_err, const string& method="")
{
if (a.size() != b.size())
{
ts->printf(CvTS::CONSOLE, "bad size");
ts->printf(CvTS::CONSOLE, "bad size, method=%s\n", method.c_str());
ts->set_failed_test_info(CvTS::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(CvTS::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(CvTS::CONSOLE, "bad accuracy: %f\n", err);
ts->printf(CvTS::CONSOLE, "bad accuracy: %f, method=%s\n", err, method.c_str());
ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);
return false;
}

Loading…
Cancel
Save