Skip to content

O
opencv
Commit 9706079a authored by Anatoly Baksheev's avatar Anatoly Baksheev
Browse files
Options

a bit refactoring in LBP face detection on GPU

parent bb516265
Hide whitespace changes
Inline Side-by-side
Showing with 121 additions and 37 deletions
modules/gpu/include/opencv2/gpu/gpu.hpp
View file @ 9706079a
......@@ -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 //////////////////////////////////////////
......
modules/gpu/src/cascadeclassifier.cpp
View file @ 9706079a
......@@ -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;
}
......
modules/gpu/test/test_objdetect.cpp
View file @ 9706079a
......@@ -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<int>(0)
));
INSTANTIATE_TEST_CASE_P(GPU_ObjDetect, LBP_Read_classifier,
testing::Combine(ALL_DEVICES, testing::Values<int>(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<cv::Rect>::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<cv::Rect>();
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<int>(0)
));
INSTANTIATE_TEST_CASE_P(GPU_ObjDetect, LBP_classify,
testing::Combine(ALL_DEVICES, testing::Values<int>(0)));
} // namespace
modules/objdetect/src/cascadedetect.cpp
View file @ 9706079a
......@@ -44,6 +44,73 @@
#include "cascadedetect.hpp"
#include <string>
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<unsigned char>(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<Rect>& candidates,
vector<int>& levels, vector<double>& 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
......
Markdown is supported
0% or
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment