From 4dfbf99dd5b5fe97128683f32e248f9cd99b46c7 Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Wed, 12 Jan 2011 09:30:08 +0000 Subject: [PATCH] cosmetic changes in gpu module, decreased matchTemplate test running time --- modules/gpu/include/opencv2/gpu/gpu.hpp | 4 +-- modules/gpu/src/cuda/match_template.cu | 20 +++++------ modules/gpu/src/hog.cpp | 10 +++--- samples/gpu/hog.cpp | 4 +-- tests/gpu/src/match_template.cpp | 48 ++++++++++++++++--------- 5 files changed, 50 insertions(+), 36 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 957299b81..70a34eee6 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1182,8 +1182,8 @@ namespace cv void setSVMDetector(const vector& detector); static vector getDefaultPeopleDetector(); - static vector getPeopleDetector_48x96(); - static vector getPeopleDetector_64x128(); + static vector getPeopleDetector48x96(); + static vector getPeopleDetector64x128(); void detect(const GpuMat& img, vector& found_locations, double hit_threshold=0, Size win_stride=Size(), diff --git a/modules/gpu/src/cuda/match_template.cu b/modules/gpu/src/cuda/match_template.cu index 7f47db830..0cdf4d1e3 100644 --- a/modules/gpu/src/cuda/match_template.cu +++ b/modules/gpu/src/cuda/match_template.cu @@ -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 diff --git a/modules/gpu/src/hog.cpp b/modules/gpu/src/hog.cpp index 9a77ad164..56f05d7ca 100644 --- a/modules/gpu/src/hog.cpp +++ b/modules/gpu/src/hog.cpp @@ -55,8 +55,8 @@ void cv::gpu::HOGDescriptor::detectMultiScale(const GpuMat&, vector&, doub void cv::gpu::HOGDescriptor::computeBlockHistograms(const GpuMat&) { throw_nogpu(); } void cv::gpu::HOGDescriptor::getDescriptors(const GpuMat&, Size, GpuMat&, int) { throw_nogpu(); } std::vector cv::gpu::HOGDescriptor::getDefaultPeopleDetector() { throw_nogpu(); return std::vector(); } -std::vector cv::gpu::HOGDescriptor::getPeopleDetector_48x96() { throw_nogpu(); return std::vector(); } -std::vector cv::gpu::HOGDescriptor::getPeopleDetector_64x128() { throw_nogpu(); return std::vector(); } +std::vector cv::gpu::HOGDescriptor::getPeopleDetector48x96() { throw_nogpu(); return std::vector(); } +std::vector cv::gpu::HOGDescriptor::getPeopleDetector64x128() { throw_nogpu(); return std::vector(); } #else @@ -352,11 +352,11 @@ cv::Size cv::gpu::HOGDescriptor::numPartsWithin(cv::Size size, cv::Size part_siz std::vector cv::gpu::HOGDescriptor::getDefaultPeopleDetector() { - return getPeopleDetector_64x128(); + return getPeopleDetector64x128(); } -std::vector cv::gpu::HOGDescriptor::getPeopleDetector_48x96() +std::vector 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 cv::gpu::HOGDescriptor::getPeopleDetector_48x96() -std::vector cv::gpu::HOGDescriptor::getPeopleDetector_64x128() +std::vector cv::gpu::HOGDescriptor::getPeopleDetector64x128() { static const float detector[] = { 0.05359386f, -0.14721455f, -0.05532170f, 0.05077307f, diff --git a/samples/gpu/hog.cpp b/samples/gpu/hog.cpp index d0ff6b3cc..81d61b0e0 100644 --- a/samples/gpu/hog.cpp +++ b/samples/gpu/hog.cpp @@ -239,9 +239,9 @@ void App::run() // Create HOG descriptors and detectors here vector 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, diff --git a/tests/gpu/src/match_template.cpp b/tests/gpu/src/match_template.cpp index 9979a54b0..2adafdc30 100644 --- a/tests/gpu/src/match_template.cpp +++ b/tests/gpu/src/match_template.cpp @@ -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(i, j); + // float b_ = b.at(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; }