From 9706079ace2ae6427533d028a1fab21ca13c1cd0 Mon Sep 17 00:00:00 2001 From: Anatoly Baksheev Date: Fri, 13 Jul 2012 15:47:09 +0000 Subject: [PATCH] a bit refactoring in LBP face detection on GPU --- modules/gpu/include/opencv2/gpu/gpu.hpp | 4 +- modules/gpu/src/cascadeclassifier.cpp | 43 +++++++------- modules/gpu/test/test_objdetect.cpp | 35 +++++++----- modules/objdetect/src/cascadedetect.cpp | 76 ++++++++++++++++++++++++- 4 files changed, 121 insertions(+), 37 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index fd72d9e9d..d30a7ec7d 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1441,7 +1441,7 @@ public: Size getClassifierSize() const; private: bool read(const FileNode &root); - void initializeBuffers(cv::Size frame); + void allocateBuffers(cv::Size frame = cv::Size()); static const stage stageType = BOOST; static const feature featureType = LBP; @@ -1463,6 +1463,8 @@ private: GpuMat integral; GpuMat integralBuffer; GpuMat resuzeBuffer; + + GpuMat candidates; }; ////////////////////////////////// SURF ////////////////////////////////////////// diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp index 37981302d..b914bbcb1 100644 --- a/modules/gpu/src/cascadeclassifier.cpp +++ b/modules/gpu/src/cascadeclassifier.cpp @@ -75,14 +75,14 @@ double /*scaleFactor*/, int /*minNeighbors*/, cv::Size /*maxObjectSize*/){ throw #else -cv::gpu::CascadeClassifier_GPU_LBP::CascadeClassifier_GPU_LBP(cv::Size detectionFrameSize) -{ - if (detectionFrameSize != cv::Size()) - initializeBuffers(detectionFrameSize); -} +cv::gpu::CascadeClassifier_GPU_LBP::CascadeClassifier_GPU_LBP(cv::Size detectionFrameSize) { allocateBuffers(detectionFrameSize); } +cv::gpu::CascadeClassifier_GPU_LBP::~CascadeClassifier_GPU_LBP(){} -void cv::gpu::CascadeClassifier_GPU_LBP::initializeBuffers(cv::Size frame) +void cv::gpu::CascadeClassifier_GPU_LBP::allocateBuffers(cv::Size frame) { + if (frame == cv::Size()) + return; + if (resuzeBuffer.empty() || frame.width > resuzeBuffer.cols || frame.height > resuzeBuffer.rows) { resuzeBuffer.create(frame, CV_8UC1); @@ -98,10 +98,12 @@ void cv::gpu::CascadeClassifier_GPU_LBP::initializeBuffers(cv::Size frame) Ncv32u bufSize; ncvSafeCall( nppiStIntegralGetSize_8u32u(roiSize, &bufSize, prop) ); integralBuffer.create(1, bufSize, CV_8UC1); + + candidates.create(1 , frame.width >> 1, CV_32SC4); } } -cv::gpu::CascadeClassifier_GPU_LBP::~CascadeClassifier_GPU_LBP(){} + void cv::gpu::CascadeClassifier_GPU_LBP::preallocateIntegralBuffer(cv::Size desired) { @@ -335,7 +337,8 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp objects.reshape(4, 1); else objects.create(1 , image.cols >> 4, CV_32SC4); - GpuMat candidates(1 , image.cols >> 1, CV_32SC4); + + candidates.create(1 , image.cols >> 1, CV_32SC4); // GpuMat candidates(1 , defaultObjSearchNum, CV_32SC4); // used for debug // candidates.setTo(cv::Scalar::all(0)); @@ -343,13 +346,12 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp if (maxObjectSize == cv::Size()) maxObjectSize = image.size(); - initializeBuffers(image.size()); - - unsigned int* classified = new unsigned int[1]; - *classified = 0; + allocateBuffers(image.size()); + + unsigned int classified = 0; unsigned int* dclassified; cudaMalloc(&dclassified, sizeof(int)); - cudaMemcpy(dclassified, classified, sizeof(int), cudaMemcpyHostToDevice); + cudaMemcpy(dclassified, &classified, sizeof(int), cudaMemcpyHostToDevice); int step = 2; // cv::gpu::device::lbp::bindIntegral(integral); @@ -370,8 +372,8 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp // if( windowSize.width < minObjectSize.width || windowSize.height < minObjectSize.height ) // continue; - GpuMat scaledImg(resuzeBuffer, cv::Rect(0, 0, scaledImageSize.width, scaledImageSize.height)); - GpuMat scaledIntegral(integral, cv::Rect(0, 0, scaledImageSize.width + 1, scaledImageSize.height + 1)); + GpuMat scaledImg = resuzeBuffer(cv::Rect(0, 0, scaledImageSize.width, scaledImageSize.height)); + GpuMat scaledIntegral = integral(cv::Rect(0, 0, scaledImageSize.width + 1, scaledImageSize.height + 1)); GpuMat currBuff = integralBuffer; cv::gpu::resize(image, scaledImg, scaledImageSize, 0, 0, CV_INTER_LINEAR); @@ -391,12 +393,13 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp // cv::gpu::device::lbp::unbindIntegral(); if (groupThreshold <= 0 || objects.empty()) return 0; - cudaMemcpy(classified, dclassified, sizeof(int), cudaMemcpyDeviceToHost); - cv::gpu::device::lbp::connectedConmonents(candidates, *classified, objects, groupThreshold, grouping_eps, dclassified); - cudaMemcpy(classified, dclassified, sizeof(int), cudaMemcpyDeviceToHost); + cudaMemcpy(&classified, dclassified, sizeof(int), cudaMemcpyDeviceToHost); + cv::gpu::device::lbp::connectedConmonents(candidates, classified, objects, groupThreshold, grouping_eps, dclassified); + cudaMemcpy(&classified, dclassified, sizeof(int), cudaMemcpyDeviceToHost); cudaSafeCall( cudaDeviceSynchronize() ); - step = *classified; - delete[] classified; + + step = classified; + cudaFree(dclassified); return step; } diff --git a/modules/gpu/test/test_objdetect.cpp b/modules/gpu/test/test_objdetect.cpp index fdd9454c8..fbd45bcc3 100644 --- a/modules/gpu/test/test_objdetect.cpp +++ b/modules/gpu/test/test_objdetect.cpp @@ -285,6 +285,10 @@ TEST_P(HOG, GetDescriptors) INSTANTIATE_TEST_CASE_P(GPU_ObjDetect, HOG, ALL_DEVICES); + +////////////////////////////////////////////////////////////////////////////////////////// +/// LBP classifier + PARAM_TEST_CASE(LBP_Read_classifier, cv::gpu::DeviceInfo, int) { cv::gpu::DeviceInfo devInfo; @@ -303,10 +307,9 @@ TEST_P(LBP_Read_classifier, Accuracy) ASSERT_TRUE(classifier.load(classifierXmlPath)); } -INSTANTIATE_TEST_CASE_P(GPU_ObjDetect, LBP_Read_classifier, testing::Combine( - ALL_DEVICES, - testing::Values(0) - )); +INSTANTIATE_TEST_CASE_P(GPU_ObjDetect, LBP_Read_classifier, + testing::Combine(ALL_DEVICES, testing::Values(0))); + PARAM_TEST_CASE(LBP_classify, cv::gpu::DeviceInfo, int) { @@ -328,7 +331,7 @@ TEST_P(LBP_classify, Accuracy) ASSERT_FALSE(cpuClassifier.empty()); cv::Mat image = cv::imread(imagePath); - image = image.colRange(0, image.cols / 2); + image = image.colRange(0, image.cols/2); cv::Mat grey; cvtColor(image, grey, CV_BGR2GRAY); ASSERT_FALSE(image.empty()); @@ -339,27 +342,29 @@ TEST_P(LBP_classify, Accuracy) std::vector::iterator it = rects.begin(); for (; it != rects.end(); ++it) - cv::rectangle(markedImage, *it, cv::Scalar(255, 0, 0, 255)); + cv::rectangle(markedImage, *it, CV_RGB(0, 0, 255)); cv::gpu::CascadeClassifier_GPU_LBP gpuClassifier; ASSERT_TRUE(gpuClassifier.load(classifierXmlPath)); + cv::gpu::GpuMat gpu_rects; cv::gpu::GpuMat tested(grey); int count = gpuClassifier.detectMultiScale(tested, gpu_rects); - cv::Mat gpu_f(gpu_rects); - int* gpu_faces = (int*)gpu_f.ptr(); + cv::Mat downloaded(gpu_rects); + const cv::Rect* faces = downloaded.ptr(); for (int i = 0; i < count; i++) { - cv::Rect r(gpu_faces[i * 4],gpu_faces[i * 4 + 1],gpu_faces[i * 4 + 2],gpu_faces[i * 4 + 3]); - std::cout << gpu_faces[i * 4]<< " " << gpu_faces[i * 4 + 1] << " " << gpu_faces[i * 4 + 2] << " " << gpu_faces[i * 4 + 3] << std::endl; - cv::rectangle(markedImage, r , cv::Scalar(0, 0, 255, 255)); + cv::Rect r = faces[i]; + + std::cout << r.x << " " << r.y << " " << r.width << " " << r.height << std::endl; + cv::rectangle(markedImage, r , CV_RGB(255, 0, 0)); } + + cv::imshow("Res", markedImage); cv::waitKey(); } -INSTANTIATE_TEST_CASE_P(GPU_ObjDetect, LBP_classify, testing::Combine( - ALL_DEVICES, - testing::Values(0) - )); +INSTANTIATE_TEST_CASE_P(GPU_ObjDetect, LBP_classify, + testing::Combine(ALL_DEVICES, testing::Values(0))); } // namespace diff --git a/modules/objdetect/src/cascadedetect.cpp b/modules/objdetect/src/cascadedetect.cpp index 91b8ba55d..e7bcac071 100644 --- a/modules/objdetect/src/cascadedetect.cpp +++ b/modules/objdetect/src/cascadedetect.cpp @@ -44,6 +44,73 @@ #include "cascadedetect.hpp" +#include + + +struct Logger +{ + enum { STADIES_NUM = 20 }; + + int gid; + cv::Mat mask; + cv::Size sz0; + int step; + + + Logger() : gid (0), step(2) {} + void setImage(const cv::Mat& image) + { + if (gid == 0) + sz0 = image.size(); + + mask.create(image.rows, image.cols * (STADIES_NUM + 1) + STADIES_NUM, CV_8UC1); + mask = cv::Scalar(0); + cv::Mat roi = mask(cv::Rect(cv::Point(0,0), image.size())); + image.copyTo(roi); + + printf("%d) Size = (%d, %d)\n", gid, image.cols, image.rows); + + for(int i = 0; i < STADIES_NUM; ++i) + { + int x = image.cols + i * (image.cols + 1); + cv::line(mask, cv::Point(x, 0), cv::Point(x, mask.rows-1), cv::Scalar(255)); + } + + if (sz0.width/image.cols > 2 && sz0.height/image.rows > 2) + step = 1; + } + + void setPoint(const cv::Point& p, int passed_stadies) + { + int cols = mask.cols / (STADIES_NUM + 1); + + passed_stadies = -passed_stadies; + passed_stadies = (passed_stadies == -1) ? STADIES_NUM : passed_stadies; + + unsigned char* ptr = mask.ptr(p.y) + cols + 1 + p.x; + for(int i = 0; i < passed_stadies; ++i, ptr += cols + 1) + { + *ptr = 255; + + if (step == 2) + { + ptr[1] = 255; + ptr[mask.step] = 255; + ptr[mask.step + 1] = 255; + } + } + }; + + void write() + { + char buf[4096]; + sprintf(buf, "%04d.png", gid++); + cv::imwrite(buf, mask); + } + +} logger; + + namespace cv { @@ -910,6 +977,8 @@ struct CascadeClassifierInvoker double gypWeight; int result = classifier->runAt(evaluator, Point(x, y), gypWeight); + + logger.setPoint(Point(x, y), result); if( rejectLevels ) { if( result == 1 ) @@ -942,6 +1011,7 @@ struct CascadeClassifierInvoker struct getRect { Rect operator ()(const CvAvgComp& e) const { return e.rect; } }; + bool CascadeClassifier::detectSingleScale( const Mat& image, int stripCount, Size processingRectSize, int stripSize, int yStep, double factor, vector& candidates, vector& levels, vector& weights, bool outputRejectLevels ) @@ -949,6 +1019,9 @@ bool CascadeClassifier::detectSingleScale( const Mat& image, int stripCount, Siz if( !featureEvaluator->setImage( image, data.origWinSize ) ) return false; + logger.setImage(image); + + Mat currentMask; if (!maskGenerator.empty()) { currentMask=maskGenerator->generateMask(image); @@ -971,7 +1044,8 @@ bool CascadeClassifier::detectSingleScale( const Mat& image, int stripCount, Siz } candidates.insert( candidates.end(), concurrentCandidates.begin(), concurrentCandidates.end() ); - return true; + logger.write(); + return true; } bool CascadeClassifier::isOldFormatCascade() const