void setSVMDetector(const vector<float>& detector);\r
bool checkDetectorSize() const;\r
\r
+ void computeBlockHistograms(const GpuMat& img);\r
void detect(const GpuMat& img, vector<Point>& found_locations, double hit_threshold=0, \r
Size win_stride=Size(), Size padding=Size());\r
void detectMultiScale(const GpuMat& img, vector<Rect>& found_locations, \r
double hit_threshold=0, Size win_stride=Size(), Size padding=Size(),\r
double scale0=1.05, int group_threshold=2);\r
-\r
- ////TODO: test it\r
- //void getDescriptors(const GpuMat& img, Size win_stride, vector<GpuMat>& descriptors)\r
+ void getDescriptors(const GpuMat& img, Size win_stride, GpuMat& descriptors);\r
\r
Size win_size;\r
Size block_size;\r
static int numPartsWithin(int size, int part_size, int stride);\r
static Size numPartsWithin(Size size, Size part_size, Size stride);\r
\r
- void computeBlockHistograms(const GpuMat& img); \r
void computeGradient(const GpuMat& img, GpuMat& grad, GpuMat& qangle);\r
\r
GpuMat grad, qangle;\r
}\r
\r
\r
-// We only support win_stride_x == block_stride_x, win_stride_y == block_stride_y\r
void classify_hists(int win_height, int win_width, int block_stride_y, int block_stride_x, \r
- int win_stride_y, int win_stride_x,\r
- int height, int width, float* block_hists, float* coefs, \r
- float free_coef, float threshold, unsigned char* labels)\r
+ int win_stride_y, int win_stride_x, int height, int width, float* block_hists, \r
+ float* coefs, float free_coef, float threshold, unsigned char* labels)\r
{ \r
const int nthreads = 256;\r
const int nblocks = 1;\r
cudaSafeCall(cudaThreadSynchronize());\r
}\r
\r
+//----------------------------------------------------------------------------\r
+// Extract descriptors\r
+\r
+\r
+template <int nthreads>\r
+__global__ void extract_descriptors_kernel(const int img_win_width, const int img_block_width, \r
+ const int win_block_stride_x, const int win_block_stride_y,\r
+ const float* block_hists, PtrElemStepf descriptors)\r
+{\r
+ // Get left top corner of the window in src\r
+ const float* hist = block_hists + (blockIdx.y * win_block_stride_y * img_block_width + \r
+ blockIdx.x * win_block_stride_x) * cblock_hist_size;\r
+\r
+ // Get left top corner of the window in dst\r
+ float* descriptor = descriptors.ptr(blockIdx.y * gridDim.x + blockIdx.x);\r
+\r
+ // Copy elements from src to dst\r
+ for (int i = threadIdx.x; i < cdescr_size; i += nthreads)\r
+ {\r
+ int offset_y = i / cdescr_width;\r
+ int offset_x = i - offset_y * cdescr_width;\r
+ descriptor[i] = hist[offset_y * img_block_width * cblock_hist_size + offset_x];\r
+ }\r
+}\r
+\r
\r
-//------------------------------------------------------------\r
+void extract_descriptors(int win_height, int win_width, int block_stride_y, int block_stride_x, \r
+ int win_stride_y, int win_stride_x, int height, int width, float* block_hists, \r
+ DevMem2Df descriptors)\r
+{\r
+ const int nthreads = 256;\r
+\r
+ int win_block_stride_x = win_stride_x / block_stride_x;\r
+ int win_block_stride_y = win_stride_y / block_stride_y;\r
+ int img_win_width = (width - win_width + win_stride_x) / win_stride_x;\r
+ int img_win_height = (height - win_height + win_stride_y) / win_stride_y;\r
+ dim3 threads(nthreads, 1);\r
+ dim3 grid(img_win_width, img_win_height);\r
+\r
+ int img_block_width = (width - CELLS_PER_BLOCK_X * CELL_WIDTH + block_stride_x) / \r
+ block_stride_x;\r
+ extract_descriptors_kernel<nthreads><<<grid, threads>>>(\r
+ img_win_width, img_block_width, win_block_stride_x, win_block_stride_y, \r
+ block_hists, descriptors);\r
+ cudaSafeCall(cudaThreadSynchronize());\r
+}\r
+\r
+//----------------------------------------------------------------------------\r
// Gradients computation\r
\r
\r
\r
float3 dx = make_float3(sqrtf(b.x) - sqrtf(a.x), \r
sqrtf(b.y) - sqrtf(a.y), \r
- sqrtf(b.z) - sqrtf(a.z)); \r
+ sqrtf(b.z) - sqrtf(a.z)); \r
float3 dy = make_float3(0.f, 0.f, 0.f);\r
\r
if (blockIdx.y > 0 && blockIdx.y < height - 1)\r
bool cv::gpu::HOGDescriptor::checkDetectorSize() const { throw_nogpu(); return false; }\r
void cv::gpu::HOGDescriptor::setSVMDetector(const vector<float>&) { throw_nogpu(); }\r
void cv::gpu::HOGDescriptor::computeGradient(const GpuMat&, GpuMat&, GpuMat&) { throw_nogpu(); }\r
+void cv::gpu::HOGDescriptor::computeBlockHistograms(const GpuMat&) { throw_nogpu(); }\r
void cv::gpu::HOGDescriptor::detect(const GpuMat&, vector<Point>&, double, Size, Size) { throw_nogpu(); }\r
void cv::gpu::HOGDescriptor::detectMultiScale(const GpuMat&, vector<Rect>&, double, Size, Size, double, int) { throw_nogpu(); }\r
+void cv::gpu::HOGDescriptor::getDescriptors(const GpuMat&, Size, GpuMat&) { throw_nogpu(); }\r
std::vector<float> cv::gpu::HOGDescriptor::getDefaultPeopleDetector() { throw_nogpu(); return std::vector<float>(); }\r
std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector_48x96() { throw_nogpu(); return std::vector<float>(); }\r
std::vector<float> cv::gpu::HOGDescriptor::getPeopleDetector_64x128() { throw_nogpu(); return std::vector<float>(); }\r
int width, float* block_hists, float* coefs, float free_coef, \r
float threshold, unsigned char* labels);\r
\r
+void extract_descriptors(int win_height, int win_width, int block_stride_y, int block_stride_x, \r
+ int win_stride_y, int win_stride_x, int height, int width, float* block_hists, \r
+ cv::gpu::DevMem2Df descriptors);\r
+\r
void compute_gradients_8UC1(int nbins, int height, int width, const cv::gpu::DevMem2D& img, \r
float angle_scale, cv::gpu::DevMem2Df grad, cv::gpu::DevMem2D qangle);\r
void compute_gradients_8UC4(int nbins, int height, int width, const cv::gpu::DevMem2D& img, \r
}\r
\r
\r
-////TODO: test it\r
-//void cv::gpu::HOGDescriptor::getDescriptors(const GpuMat& img, Size win_stride, \r
-// vector<GpuMat>& descriptors)\r
-//{\r
-// CV_Assert(win_stride.width % block_stride.width == 0 &&\r
-// win_stride.height % block_stride.height == 0);\r
-//\r
-// computeBlockHistograms(img);\r
-//\r
-// Size blocks_per_img = numPartsWithin(img.size(), block_size, block_stride);\r
-// GpuMat hists_reshaped = block_hists.reshape(0, blocks_per_img.height);\r
-//\r
-// const int block_hist_size = getBlockHistogramSize();\r
-// Size blocks_per_win = numPartsWithin(win_size, block_size, block_stride);\r
-// Size wins_per_img = numPartsWithin(img.size(), win_size, win_stride);\r
-//\r
-// descriptors.resize(wins_per_img.area());\r
-// for (int i = 0; i < wins_per_img.height; ++i)\r
-// {\r
-// for (int j = 0; j < wins_per_img.width; ++j)\r
-// {\r
-// Range rows;\r
-// rows.start = i * (blocks_per_win.height + 1);\r
-// rows.end = rows.start + blocks_per_win.height;\r
-//\r
-// Range cols;\r
-// cols.start = j * (blocks_per_win.width + 1) * block_hist_size;\r
-// cols.end = cols.start + blocks_per_win.width * block_hist_size;\r
-//\r
-// descriptors[i * wins_per_img.width + j] = hists_reshaped(rows, cols);\r
-// }\r
-// }\r
-//}\r
+void cv::gpu::HOGDescriptor::getDescriptors(const GpuMat& img, Size win_stride, GpuMat& descriptors)\r
+{\r
+ CV_Assert(win_stride.width % block_stride.width == 0 &&\r
+ win_stride.height % block_stride.height == 0);\r
+\r
+ computeBlockHistograms(img);\r
+\r
+ const int block_hist_size = getBlockHistogramSize();\r
+ Size blocks_per_win = numPartsWithin(win_size, block_size, block_stride);\r
+ Size wins_per_img = numPartsWithin(img.size(), win_size, win_stride);\r
+\r
+ descriptors.create(wins_per_img.area(), blocks_per_win.area() * block_hist_size, CV_32F);\r
+\r
+ hog::extract_descriptors(win_size.height, win_size.width, block_stride.height, block_stride.width, \r
+ win_stride.height, win_stride.width, img.rows, img.cols, block_hists.ptr<float>(), \r
+ descriptors);\r
+}\r
\r
\r
void cv::gpu::HOGDescriptor::detect(const GpuMat& img, vector<Point>& hits, double hit_threshold, \r
vc >> frame;\r
}\r
else\r
+ {\r
frame = imread(settings.src);\r
+ if (frame.empty())\r
+ throw exception(string("Can't open image file: " + settings.src).c_str());\r
+ }\r
\r
Mat img_aux, img, img_to_show;\r
gpu::GpuMat gpu_img;\r
ts->set_failed_test_info(err); \\r
return; }\r
\r
-struct CV_GpuHogTest : public CvTest \r
+struct CV_GpuHogDetectionTest: public CvTest \r
{\r
- CV_GpuHogTest() : CvTest( "GPU-HOG", "HOGDescriptor" ) {}\r
+ CV_GpuHogDetectionTest(): CvTest( "GPU-HOG-detect", "HOGDescriptorDetection" ) {}\r
+\r
+ void run(int) \r
+ { \r
+ try \r
+ {\r
+ cv::Mat img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/road.png");\r
+ CHECK(!img_rgb.empty(), CvTS::FAIL_MISSING_TEST_DATA);\r
+\r
+#ifdef DUMP\r
+ f.open((std::string(ts->get_data_path()) + "hog/expected_output.bin").c_str(), std::ios_base::binary);\r
+ CHECK(f.is_open(), CvTS::FAIL_GENERIC); \r
+#else\r
+ f.open((std::string(ts->get_data_path()) + "hog/expected_output.bin").c_str(), std::ios_base::binary);\r
+ CHECK(f.is_open(), CvTS::FAIL_MISSING_TEST_DATA); \r
+#endif\r
+ \r
+ // Test on color image\r
+ cv::Mat img;\r
+ cv::cvtColor(img_rgb, img, CV_BGR2BGRA);\r
+ test(img);\r
+\r
+ // Test on gray image\r
+ cv::cvtColor(img_rgb, img, CV_BGR2GRAY);\r
+ test(img);\r
+\r
+ f.close();\r
+ }\r
+ catch (const cv::Exception& e)\r
+ {\r
+ f.close();\r
+ if (!check_and_treat_gpu_exception(e, ts)) throw;\r
+ return;\r
+ }\r
+ }\r
\r
#ifdef DUMP\r
void dump(const cv::Mat& block_hists, const std::vector<cv::Point>& locations) \r
#endif\r
}\r
\r
+#ifdef DUMP\r
+ std::ofstream f;\r
+#else\r
+ std::ifstream f;\r
+#endif\r
\r
- void run(int) \r
- { \r
+} gpu_hog_detection_test;\r
+\r
+\r
+struct CV_GpuHogGetDescriptorsTest: public CvTest \r
+{\r
+ CV_GpuHogGetDescriptorsTest(): CvTest("GPU-HOG-getDescriptors", "HOGDescriptorGetDescriptors") {}\r
+\r
+ void run(int)\r
+ {\r
try \r
{\r
- cv::Mat img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/road.png");\r
+ // Load image (e.g. train data, composed from windows)\r
+ cv::Mat img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/train_data.png");\r
CHECK(!img_rgb.empty(), CvTS::FAIL_MISSING_TEST_DATA);\r
\r
-#ifdef DUMP\r
- f.open((std::string(ts->get_data_path()) + "hog/expected_output.bin").c_str(), std::ios_base::binary);\r
- CHECK(f.is_open(), CvTS::FAIL_GENERIC); \r
-#else\r
- f.open((std::string(ts->get_data_path()) + "hog/expected_output.bin").c_str(), std::ios_base::binary);\r
- CHECK(f.is_open(), CvTS::FAIL_MISSING_TEST_DATA); \r
-#endif\r
- \r
- // Test on color image\r
+ // Convert to C4\r
cv::Mat img;\r
cv::cvtColor(img_rgb, img, CV_BGR2BGRA);\r
- test(img);\r
+ cv::gpu::GpuMat d_img(img);\r
\r
- // Test on gray image\r
- cv::cvtColor(img_rgb, img, CV_BGR2GRAY);\r
- test(img);\r
+ cv::Size win_size(64, 128);\r
+ cv::gpu::HOGDescriptor hog(win_size);\r
\r
- f.close();\r
+ // Convert train images into feature vectors (train table)\r
+ cv::gpu::GpuMat descriptors;\r
+ hog.getDescriptors(d_img, win_size, descriptors);\r
+\r
+ // Check size of the result train table\r
+ wins_per_img_x = 3;\r
+ wins_per_img_y = 2;\r
+ blocks_per_win_x = 7;\r
+ blocks_per_win_y = 15;\r
+ block_hist_size = 36;\r
+ cv::Size descr_size_expected = cv::Size(blocks_per_win_x * blocks_per_win_y * block_hist_size, \r
+ wins_per_img_x * wins_per_img_y); \r
+ CHECK(descriptors.size() == descr_size_expected, CvTS::FAIL_INVALID_OUTPUT);\r
+\r
+ /* Now we want to extract the same feature vectors, but from single images. NOTE: results will \r
+ be defferent, due to border values interpolation. Using of many small images is slower, however we \r
+ wont't call getDescriptors and will use computeBlockHistograms instead of. computeBlockHistograms \r
+ works good, it can be checked in the gpu_hog sample */\r
+\r
+ img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/positive1.png");\r
+ CHECK(!img_rgb.empty(), CvTS::FAIL_MISSING_TEST_DATA);\r
+ cv::cvtColor(img_rgb, img, CV_BGR2BGRA);\r
+ hog.computeBlockHistograms(cv::gpu::GpuMat(img));\r
+ // Everything is fine with interpolation for left top subimage\r
+ CHECK(cv::norm(hog.block_hists, descriptors.rowRange(0, 1)) == 0.f, CvTS::FAIL_INVALID_OUTPUT);\r
+\r
+ img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/positive2.png");\r
+ CHECK(!img_rgb.empty(), CvTS::FAIL_MISSING_TEST_DATA);\r
+ cv::cvtColor(img_rgb, img, CV_BGR2BGRA);\r
+ hog.computeBlockHistograms(cv::gpu::GpuMat(img));\r
+ compare_inner_parts(hog.block_hists, descriptors.rowRange(1, 2));\r
+\r
+ img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/negative1.png");\r
+ CHECK(!img_rgb.empty(), CvTS::FAIL_MISSING_TEST_DATA);\r
+ cv::cvtColor(img_rgb, img, CV_BGR2BGRA);\r
+ hog.computeBlockHistograms(cv::gpu::GpuMat(img));\r
+ compare_inner_parts(hog.block_hists, descriptors.rowRange(2, 3));\r
+\r
+ img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/negative2.png");\r
+ CHECK(!img_rgb.empty(), CvTS::FAIL_MISSING_TEST_DATA);\r
+ cv::cvtColor(img_rgb, img, CV_BGR2BGRA);\r
+ hog.computeBlockHistograms(cv::gpu::GpuMat(img));\r
+ compare_inner_parts(hog.block_hists, descriptors.rowRange(3, 4));\r
+\r
+ img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/positive3.png");\r
+ CHECK(!img_rgb.empty(), CvTS::FAIL_MISSING_TEST_DATA);\r
+ cv::cvtColor(img_rgb, img, CV_BGR2BGRA);\r
+ hog.computeBlockHistograms(cv::gpu::GpuMat(img));\r
+ compare_inner_parts(hog.block_hists, descriptors.rowRange(4, 5));\r
+\r
+ img_rgb = cv::imread(std::string(ts->get_data_path()) + "hog/negative3.png");\r
+ CHECK(!img_rgb.empty(), CvTS::FAIL_MISSING_TEST_DATA);\r
+ cv::cvtColor(img_rgb, img, CV_BGR2BGRA);\r
+ hog.computeBlockHistograms(cv::gpu::GpuMat(img));\r
+ compare_inner_parts(hog.block_hists, descriptors.rowRange(5, 6));\r
}\r
catch (const cv::Exception& e)\r
{\r
- f.close();\r
if (!check_and_treat_gpu_exception(e, ts)) throw;\r
return;\r
}\r
}\r
\r
-#ifdef DUMP\r
- std::ofstream f;\r
-#else\r
- std::ifstream f;\r
-#endif\r
+ // Does not compare border value, as interpolation leads to delta\r
+ void compare_inner_parts(cv::Mat d1, cv::Mat d2)\r
+ {\r
+ for (int i = 1; i < blocks_per_win_y - 1; ++i)\r
+ for (int j = 1; j < blocks_per_win_x - 1; ++j)\r
+ for (int k = 0; k < block_hist_size; ++k)\r
+ {\r
+ float a = d1.at<float>(0, (i * blocks_per_win_x + j) * block_hist_size);\r
+ float b = d2.at<float>(0, (i * blocks_per_win_x + j) * block_hist_size);\r
+ CHECK(a == b, CvTS::FAIL_INVALID_OUTPUT)\r
+ }\r
+ }\r
+\r
+ int wins_per_img_x;\r
+ int wins_per_img_y;\r
+ int blocks_per_win_x;\r
+ int blocks_per_win_y;\r
+ int block_hist_size;\r
+} gpu_hog_get_descriptors_test;\r
\r
-} gpu_hog_test;\r
projects / profile / ivi / opencv.git / commitdiff