From fada88e100406e71e6d34cacf601e1bd7e663e07 Mon Sep 17 00:00:00 2001 From: Alexey Spizhevoy Date: Mon, 15 Nov 2010 12:51:30 +0000 Subject: [PATCH] fixed some bugs, added tests for gpu::HOGDescriptor --- modules/gpu/include/opencv2/gpu/gpu.hpp | 4 +- modules/gpu/src/cuda/hog.cu | 1 + modules/gpu/src/hog.cpp | 5 - modules/gpu/src/matrix_operations.cpp | 2 +- tests/gpu/src/hog.cpp | 210 ++++++++++++++++++++++++++++++++ 5 files changed, 213 insertions(+), 9 deletions(-) create mode 100644 tests/gpu/src/hog.cpp diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 50e1f90..0fdec1e 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -282,7 +282,7 @@ namespace cv operator GpuMat() const; //returns if host memory can be mapperd to gpu address space; - static bool can_device_map_to_host(); + static bool canMapHostMemory(); // Please see cv::Mat for descriptions bool isContinuous() const; @@ -1001,8 +1001,6 @@ namespace cv double hit_threshold=0, Size win_stride=Size(), Size padding=Size(), double scale0=1.05, int group_threshold=2); - GpuMat getLastBlockHists() const; - Size win_size; Size block_size; Size block_stride; diff --git a/modules/gpu/src/cuda/hog.cu b/modules/gpu/src/cuda/hog.cu index ccf974f..371ce7f 100644 --- a/modules/gpu/src/cuda/hog.cu +++ b/modules/gpu/src/cuda/hog.cu @@ -142,6 +142,7 @@ __global__ void compute_hists_kernel_many_blocks(const int img_block_width, cons const float* grad_ptr = grad.ptr(offset_y) + offset_x * 2; const unsigned char* qangle_ptr = qangle.ptr(offset_y) + offset_x * 2; + // 12 means that 12 pixels affect on block's cell (in one row) if (cell_thread_x < 12) { float* hist = hists + 12 * (cell_y * blockDim.z * CELLS_PER_BLOCK_Y + diff --git a/modules/gpu/src/hog.cpp b/modules/gpu/src/hog.cpp index cfb79a5..88e9721 100644 --- a/modules/gpu/src/hog.cpp +++ b/modules/gpu/src/hog.cpp @@ -295,11 +295,6 @@ void cv::gpu::HOGDescriptor::detectMultiScale(const GpuMat& img, vector& f } -cv::gpu::GpuMat cv::gpu::HOGDescriptor::getLastBlockHists() const { - return block_hists; -} - - int cv::gpu::HOGDescriptor::numPartsWithin(int size, int part_size, int stride) { return (size - part_size + stride) / stride; diff --git a/modules/gpu/src/matrix_operations.cpp b/modules/gpu/src/matrix_operations.cpp index 153c524..cebb6a3 100644 --- a/modules/gpu/src/matrix_operations.cpp +++ b/modules/gpu/src/matrix_operations.cpp @@ -68,7 +68,7 @@ namespace cv void GpuMat::release() { throw_nogpu(); } void CudaMem::create(int /*_rows*/, int /*_cols*/, int /*_type*/, int /*type_alloc*/) { throw_nogpu(); } - bool CudaMem::can_device_map_to_host() { throw_nogpu(); return false; } + bool CudaMem::canMapHostMemory() { throw_nogpu(); return false; } void CudaMem::release() { throw_nogpu(); } GpuMat CudaMem::createGpuMatHeader () const { throw_nogpu(); return GpuMat(); } } diff --git a/tests/gpu/src/hog.cpp b/tests/gpu/src/hog.cpp new file mode 100644 index 0000000..e3f2719 --- /dev/null +++ b/tests/gpu/src/hog.cpp @@ -0,0 +1,210 @@ +/*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 "gputest.hpp" +#include + +using namespace std; + +//#define DUMP + +#define CHECK(pred, err) if (!(pred)) { \ + ts->printf(CvTS::LOG, "Fail: \"%s\" at line: %d\n", #pred, __LINE__); \ + ts->set_failed_test_info(err); \ + return; } + +struct CV_GpuHogTest : public CvTest +{ + CV_GpuHogTest() : CvTest( "GPU-HOG", "HOGDescriptor" ) {} + +#ifdef DUMP + void dump(const cv::Mat& block_hists, const std::vector& locations) + { + f.write((char*)&block_hists.rows, sizeof(block_hists.rows)); + f.write((char*)&block_hists.cols, sizeof(block_hists.cols)); + for (int i = 0; i < block_hists.rows; ++i) + { + for (int j = 0; j < block_hists.cols; ++j) + { + float val = block_hists.at(i, j); + f.write((char*)&val, sizeof(val)); + } + } + size_t nlocations = locations.size(); + f.write((char*)&nlocations, sizeof(nlocations)); + for (size_t i = 0; i < locations.size(); ++i) + f.write((char*)&locations[i], sizeof(locations[i])); + } +#else + void compare(const cv::Mat& block_hists, const std::vector& locations) + { + int rows, cols; + size_t nlocations; + + f.read((char*)&rows, sizeof(rows)); + f.read((char*)&cols, sizeof(cols)); + CHECK(rows == block_hists.rows, CvTS::FAIL_INVALID_OUTPUT); + CHECK(cols == block_hists.cols, CvTS::FAIL_INVALID_OUTPUT); + 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(i, j)) < 1e-3f, CvTS::FAIL_INVALID_OUTPUT); + } + } + f.read((char*)&nlocations, sizeof(nlocations)); + CHECK(nlocations == locations.size(), CvTS::FAIL_INVALID_OUTPUT); + for (size_t i = 0; i < nlocations; ++i) + { + cv::Point location; + f.read((char*)&location, sizeof(location)); + CHECK(location == locations[i], CvTS::FAIL_INVALID_OUTPUT); + } + } +#endif + + void test(const cv::Mat& img) + { + cv::gpu::GpuMat d_img(img); + + cv::gpu::HOGDescriptor hog; + hog.setSVMDetector(cv::gpu::HOGDescriptor::getDefaultPeopleDetector()); + + std::vector locations; + + // Test detect + hog.detect(d_img, locations, 0); + +#ifdef DUMP + dump(hog.block_hists, locations); +#else + compare(hog.block_hists, locations); +#endif + + // Test detect on smaller image + cv::gpu::GpuMat d_img2; + cv::gpu::resize(d_img, d_img2, cv::Size(d_img.cols / 2, d_img.rows / 2)); + hog.detect(d_img2, locations, 0); + +#ifdef DUMP + dump(hog.block_hists, locations); +#else + compare(hog.block_hists, locations); +#endif + + // Test detect on greater image + cv::gpu::resize(d_img, d_img2, cv::Size(d_img.cols * 2, d_img.rows * 2)); + hog.detect(d_img2, locations, 0); + +#ifdef DUMP + dump(hog.block_hists, locations); +#else + compare(hog.block_hists, locations); +#endif + + // Test detectMultiScale + std::vector rects; + size_t nrects; + hog.detectMultiScale(d_img, rects, 0, cv::Size(8, 8), cv::Size(), 1.05, 2); + +#ifdef DUMP + nrects = rects.size(); + f.write((char*)&nrects, sizeof(nrects)); + for (size_t i = 0; i < rects.size(); ++i) + f.write((char*)&rects[i], sizeof(rects[i])); + dump(hog.block_hists, std::vector()); +#else + f.read((char*)&nrects, sizeof(nrects)); + CHECK(nrects == rects.size(), CvTS::FAIL_INVALID_OUTPUT) + for (size_t i = 0; i < rects.size(); ++i) + { + cv::Rect rect; + f.read((char*)&rect, sizeof(rect)); + CHECK(rect == rects[i], CvTS::FAIL_INVALID_OUTPUT); + } + compare(hog.block_hists, std::vector()); +#endif + } + + + void run(int) + { + try + { + cv::Mat img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/road.png"); + CHECK(!img_rgb.empty(), CvTS::FAIL_MISSING_TEST_DATA); + +#ifdef DUMP + f.open((std::string(ts->get_data_path()) + "hog/expected_output.bin").c_str(), std::ios_base::binary); + CHECK(f.is_open(), CvTS::FAIL_GENERIC); +#else + f.open((std::string(ts->get_data_path()) + "hog/expected_output.bin").c_str(), std::ios_base::binary); + CHECK(f.is_open(), CvTS::FAIL_MISSING_TEST_DATA); +#endif + + // Test on color image + cv::Mat img; + cv::cvtColor(img_rgb, img, CV_BGR2BGRA); + test(img); + + // Test on gray image + cv::cvtColor(img_rgb, img, CV_BGR2GRAY); + test(img); + + f.close(); + } + catch (const cv::Exception& e) + { + f.close(); + if (!check_and_treat_gpu_exception(e, ts)) throw; + return; + } + } + +#ifdef DUMP + std::ofstream f; +#else + std::ifstream f; +#endif + +} gpu_hog_test; -- 2.7.4