From c470e15d4598cbb5927e836082f739e8903fedb4 Mon Sep 17 00:00:00 2001 From: "marina.kolpakova" Date: Wed, 12 Dec 2012 03:43:20 +0400 Subject: [PATCH] integrate speprocessing strategy --- modules/gpu/include/opencv2/gpu/gpu.hpp | 61 +++--- modules/gpu/perf/perf_softcascade.cpp | 36 ++-- modules/gpu/src/gpu_init.cpp | 7 +- modules/gpu/src/softcascade.cpp | 343 ++++++++++++++------------------ modules/gpu/test/test_softcascade.cpp | 21 +- 5 files changed, 204 insertions(+), 264 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 7003c8f..866c3de 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1529,32 +1529,37 @@ public: // ======================== GPU version for soft cascade ===================== // -// Implementation of soft (stageless) cascaded detector. -class CV_EXPORTS SCascade : public Algorithm +class CV_EXPORTS ChannelsProcessor { public: - - enum { GENERIC = 1, SEPARABLE = 2}; - class CV_EXPORTS Preprocessor + enum { - public: + GENERIC = 1 << 4, + SEPARABLE = 2 << 4 + }; - // Appends specified number of HOG first-order features integrals into given vector. - // Param frame is an input 3-channel bgr image. - // Param channels is a GPU matrix of integrals. - // Param stream is stream is a high-level CUDA stream abstraction used for asynchronous execution. - virtual void apply(InputArray frame, OutputArray channels, Stream& stream = Stream::Null()) = 0; + // Appends specified number of HOG first-order features integrals into given vector. + // Param frame is an input 3-channel bgr image. + // Param channels is a GPU matrix of optionally shrinked channels + // Param stream is stream is a high-level CUDA stream abstraction used for asynchronous execution. + virtual void apply(InputArray frame, OutputArray channels, Stream& stream = Stream::Null()) = 0; - // Creates a specific preprocessor implementation. - // Param shrinkage is a resizing factor. Resize is applied before the computing integral sum - // Param bins is a number of HOG-like channels. - // Param method is a channel computing method. - static cv::Ptr create(const int shrinkage, const int bins, const int method = GENERIC); + // Creates a specific preprocessor implementation. + // Param shrinkage is a resizing factor. Resize is applied before the computing integral sum + // Param bins is a number of HOG-like channels. + // Param flags is a channel computing extra flags. + static cv::Ptr create(const int shrinkage, const int bins, const int flags = GENERIC); + virtual ~ChannelsProcessor(); - protected: - Preprocessor(); - }; +protected: + ChannelsProcessor(); +}; + +// Implementation of soft (stageless) cascaded detector. +class CV_EXPORTS SCascade : public Algorithm +{ +public: // Representation of detectors result. struct CV_EXPORTS Detection @@ -1569,14 +1574,15 @@ public: enum {PEDESTRIAN = 0}; }; - enum { NO_REJECT = 1, DOLLAR = 2, /*PASCAL = 4,*/ DEFAULT = NO_REJECT}; + enum { NO_REJECT = 1, DOLLAR = 2, /*PASCAL = 4,*/ DEFAULT = NO_REJECT, NMS_MASK = 0xF}; // An empty cascade will be created. // Param minScale is a minimum scale relative to the original size of the image on which cascade will be applyed. // Param minScale is a maximum scale relative to the original size of the image on which cascade will be applyed. // Param scales is a number of scales from minScale to maxScale. - // Param rejfactor is used for NMS. - SCascade(const double minScale = 0.4, const double maxScale = 5., const int scales = 55, const int rejCriteria = 1); + // Param flags is an extra tuning flags. + SCascade(const double minScale = 0.4, const double maxScale = 5., const int scales = 55, + const int flags = NO_REJECT || ChannelsProcessor::GENERIC); virtual ~SCascade(); @@ -1598,13 +1604,6 @@ public: // Param stream is stream is a high-level CUDA stream abstraction used for asynchronous execution virtual void detect(InputArray image, InputArray rois, OutputArray objects, Stream& stream = Stream::Null()) const; - // Convert ROI matrix into the suitable for detect method. - // Param roi is an input matrix of the same size as the image. - // There non zero value mean that detector should be executed in this point. - // Param mask is an output mask - // Param stream is stream is a high-level CUDA stream abstraction used for asynchronous execution - virtual void genRoi(InputArray roi, OutputArray mask, Stream& stream = Stream::Null()) const; - private: struct Fields; @@ -1612,9 +1611,9 @@ private: double minScale; double maxScale; - int scales; - int rejCriteria; + + int flags; }; CV_EXPORTS bool initModule_gpu(void); diff --git a/modules/gpu/perf/perf_softcascade.cpp b/modules/gpu/perf/perf_softcascade.cpp index b4091d0..9645bee 100644 --- a/modules/gpu/perf/perf_softcascade.cpp +++ b/modules/gpu/perf/perf_softcascade.cpp @@ -71,15 +71,14 @@ RUN_GPU(SCascadeTest, detect) ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); - cv::gpu::GpuMat objectBoxes(1, 10000 * sizeof(cv::gpu::SCascade::Detection), CV_8UC1), rois(colored.size(), CV_8UC1), trois; + cv::gpu::GpuMat objectBoxes(1, 10000 * sizeof(cv::gpu::SCascade::Detection), CV_8UC1), rois(colored.size(), CV_8UC1); rois.setTo(1); - cascade.genRoi(rois, trois); - cascade.detect(colored, trois, objectBoxes); + cascade.detect(colored, rois, objectBoxes); TEST_CYCLE() { - cascade.detect(colored, trois, objectBoxes); + cascade.detect(colored, rois, objectBoxes); } SANITY_CHECK(sortDetections(objectBoxes)); @@ -142,14 +141,11 @@ RUN_GPU(SCascadeTestRoi, detectInRoi) sub.setTo(1); } - cv::gpu::GpuMat trois; - cascade.genRoi(rois, trois); - - cascade.detect(colored, trois, objectBoxes); + cascade.detect(colored, rois, objectBoxes); TEST_CYCLE() { - cascade.detect(colored, trois, objectBoxes); + cascade.detect(colored, rois, objectBoxes); } SANITY_CHECK(sortDetections(objectBoxes)); @@ -186,14 +182,11 @@ RUN_GPU(SCascadeTestRoi, detectEachRoi) cv::gpu::GpuMat sub(rois, r); sub.setTo(1); - cv::gpu::GpuMat trois; - cascade.genRoi(rois, trois); - - cascade.detect(colored, trois, objectBoxes); + cascade.detect(colored, rois, objectBoxes); TEST_CYCLE() { - cascade.detect(colored, trois, objectBoxes); + cascade.detect(colored, rois, objectBoxes); } SANITY_CHECK(sortDetections(objectBoxes)); @@ -235,15 +228,14 @@ RUN_GPU(SCascadeTest, detectOnIntegral) ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); - cv::gpu::GpuMat objectBoxes(1, 10000 * sizeof(cv::gpu::SCascade::Detection), CV_8UC1), rois(cv::Size(640, 480), CV_8UC1), trois; + cv::gpu::GpuMat objectBoxes(1, 10000 * sizeof(cv::gpu::SCascade::Detection), CV_8UC1), rois(cv::Size(640, 480), CV_8UC1); rois.setTo(1); - cascade.genRoi(rois, trois); - cascade.detect(hogluv, trois, objectBoxes); + cascade.detect(hogluv, rois, objectBoxes); TEST_CYCLE() { - cascade.detect(hogluv, trois, objectBoxes); + cascade.detect(hogluv, rois, objectBoxes); } SANITY_CHECK(sortDetections(objectBoxes)); @@ -270,18 +262,16 @@ RUN_GPU(SCascadeTest, detectStream) ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); - cv::gpu::GpuMat objectBoxes(1, 10000 * sizeof(cv::gpu::SCascade::Detection), CV_8UC1), rois(colored.size(), CV_8UC1), trois; + cv::gpu::GpuMat objectBoxes(1, 10000 * sizeof(cv::gpu::SCascade::Detection), CV_8UC1), rois(colored.size(), CV_8UC1); rois.setTo(1); cv::gpu::Stream s; - cascade.genRoi(rois, trois, s); - - cascade.detect(colored, trois, objectBoxes, s); + cascade.detect(colored, rois, objectBoxes, s); TEST_CYCLE() { - cascade.detect(colored, trois, objectBoxes, s); + cascade.detect(colored, rois, objectBoxes, s); } #ifdef HAVE_CUDA diff --git a/modules/gpu/src/gpu_init.cpp b/modules/gpu/src/gpu_init.cpp index 773a8b6..cffacb8 100644 --- a/modules/gpu/src/gpu_init.cpp +++ b/modules/gpu/src/gpu_init.cpp @@ -46,10 +46,9 @@ namespace cv { namespace gpu { CV_INIT_ALGORITHM(SCascade, "CascadeDetector.SCascade", - obj.info()->addParam(obj, "minScale", obj.minScale); - obj.info()->addParam(obj, "maxScale", obj.maxScale); - obj.info()->addParam(obj, "scales", obj.scales); - obj.info()->addParam(obj, "rejCriteria", obj.rejCriteria)); + obj.info()->addParam(obj, "minScale", obj.minScale); + obj.info()->addParam(obj, "maxScale", obj.maxScale); + obj.info()->addParam(obj, "scales", obj.scales)); bool initModule_gpu(void) { diff --git a/modules/gpu/src/softcascade.cpp b/modules/gpu/src/softcascade.cpp index 6d29a1d..9c9132b 100644 --- a/modules/gpu/src/softcascade.cpp +++ b/modules/gpu/src/softcascade.cpp @@ -41,10 +41,8 @@ //M*/ #include -#include #if !defined (HAVE_CUDA) - cv::gpu::SCascade::SCascade(const double, const double, const int, const int) { throw_nogpu(); } cv::gpu::SCascade::~SCascade() { throw_nogpu(); } @@ -53,18 +51,16 @@ bool cv::gpu::SCascade::load(const FileNode&) { throw_nogpu(); return false;} void cv::gpu::SCascade::detect(InputArray, InputArray, OutputArray, Stream&) const { throw_nogpu(); } -void cv::gpu::SCascade::genRoi(InputArray, OutputArray, Stream&) const { throw_nogpu(); } - void cv::gpu::SCascade::read(const FileNode& fn) { Algorithm::read(fn); } -cv::gpu::SCascade::Preprocessor::Preprocessor() { throw_nogpu(); } - -void cv::gpu::SCascade::Preprocessor::create(const int, const int, const int) { throw_nogpu(); } +cv::gpu::ChannelsProcessor::ChannelsProcessor() { throw_nogpu(); } + cv::gpu::ChannelsProcessor::~ChannelsProcessor() { throw_nogpu(); } +cv::Ptr cv::gpu::ChannelsProcessor::create(const int, const int, const int) +{ throw_nogpu(); return cv::Ptr(0); } #else - -#include +# include cv::gpu::device::icf::Level::Level(int idx, const Octave& oct, const float scale, const int w, const int h) : octave(idx), step(oct.stages), relScale(scale / oct.scale) @@ -96,23 +92,22 @@ namespace icf { void bgr2Luv(const PtrStepSzb& bgr, PtrStepSzb luv); void gray2hog(const PtrStepSzb& gray, PtrStepSzb mag, const int bins); - void shrink(const cv::gpu::PtrStepSzb& channels, cv::gpu::PtrStepSzb shrunk); } -namespace imgproc { - void shfl_integral_gpu_buffered(PtrStepSzb, PtrStepSz, PtrStepSz, int, cudaStream_t); +// namespace imgproc { +// void shfl_integral_gpu_buffered(PtrStepSzb, PtrStepSz, PtrStepSz, int, cudaStream_t); - template - void resize_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float fx, float fy, - PtrStepSzb dst, int interpolation, cudaStream_t stream); -} +// template +// void resize_gpu(PtrStepSzb src, PtrStepSzb srcWhole, int xoff, int yoff, float fx, float fy, +// PtrStepSzb dst, int interpolation, cudaStream_t stream); +// } }}} struct cv::gpu::SCascade::Fields { - static Fields* parseCascade(const FileNode &root, const float mins, const float maxs, const int totals) + static Fields* parseCascade(const FileNode &root, const float mins, const float maxs, const int totals, const int method) { static const char *const SC_STAGE_TYPE = "stageType"; static const char *const SC_BOOST = "BOOST"; @@ -253,9 +248,9 @@ struct cv::gpu::SCascade::Fields CV_Assert(!hleaves.empty()); Fields* fields = new Fields(mins, maxs, totals, origWidth, origHeight, shrinkage, 0, - hoctaves, hstages, hnodes, hleaves); + hoctaves, hstages, hnodes, hleaves, method); fields->voctaves = voctaves; - fields->createLevels(FRAME_HEIGHT, FRAME_WIDTH); + fields->createLevels(DEFAULT_FRAME_HEIGHT, DEFAULT_FRAME_WIDTH); return fields; } @@ -310,12 +305,6 @@ struct cv::gpu::SCascade::Fields bool update(int fh, int fw, int shr) { - if ((fh == luv.rows) && (fw == luv.cols)) return false; - - plane.create(fh * (HOG_LUV_BINS + 1), fw, CV_8UC1); - fplane.create(fh * HOG_BINS, fw, CV_32FC1); - luv.create(fh, fw, CV_8UC3); - shrunk.create(fh / shr * HOG_LUV_BINS, fw / shr, CV_8UC1); integralBuffer.create(shrunk.rows, shrunk.cols, CV_32SC1); @@ -329,17 +318,19 @@ struct cv::gpu::SCascade::Fields } Fields( const float mins, const float maxs, const int tts, const int ow, const int oh, const int shr, const int ds, - cv::Mat hoctaves, cv::Mat hstages, cv::Mat hnodes, cv::Mat hleaves) + cv::Mat hoctaves, cv::Mat hstages, cv::Mat hnodes, cv::Mat hleaves, int method) : minScale(mins), maxScale(maxs), totals(tts), origObjWidth(ow), origObjHeight(oh), shrinkage(shr), downscales(ds) { - update(FRAME_HEIGHT, FRAME_WIDTH, shr); + update(DEFAULT_FRAME_HEIGHT, DEFAULT_FRAME_WIDTH, shr); octaves.upload(hoctaves); stages.upload(hstages); nodes.upload(hnodes); leaves.upload(hleaves); + + preprocessor = ChannelsProcessor::create(shrinkage, 6, method); } - void detect(const cv::gpu::GpuMat& roi, cv::gpu::GpuMat& objects, Stream& s) const + void detect(cv::gpu::GpuMat& objects, Stream& s) const { if (s) s.enqueueMemSet(objects, 0); @@ -352,26 +343,7 @@ struct cv::gpu::SCascade::Fields = device::icf::CascadeInvoker(levels, stages, nodes, leaves); cudaStream_t stream = StreamAccessor::getStream(s); - invoker(roi, hogluv, objects, downscales, stream); - } - - void preprocess(const cv::gpu::GpuMat& colored, Stream& s) - { - if (s) - s.enqueueMemSet(plane, 0); - else - cudaMemset(plane.data, 0, plane.step * plane.rows); - - const int fw = colored.cols; - const int fh = colored.rows; - - GpuMat gray(plane, cv::Rect(0, fh * Fields::HOG_LUV_BINS, fw, fh)); - cv::gpu::cvtColor(colored, gray, CV_BGR2GRAY, s); - createHogBins(gray ,s); - - createLuvBins(colored, s); - - integrate(fh, fw, s); + invoker(mask, hogluv, objects, downscales, stream); } void suppress(GpuMat& objects, Stream& s) @@ -416,72 +388,10 @@ private: return res; } - void createHogBins(const cv::gpu::GpuMat& gray, Stream& s) - { - static const int fw = gray.cols; - static const int fh = gray.rows; - - GpuMat dfdx(fplane, cv::Rect(0, 0, fw, fh)); - GpuMat dfdy(fplane, cv::Rect(0, fh, fw, fh)); - - cv::gpu::Sobel(gray, dfdx, CV_32F, 1, 0, sobelBuf, 3, 1, BORDER_DEFAULT, -1, s); - cv::gpu::Sobel(gray, dfdy, CV_32F, 0, 1, sobelBuf, 3, 1, BORDER_DEFAULT, -1, s); - - GpuMat mag(fplane, cv::Rect(0, 2 * fh, fw, fh)); - GpuMat ang(fplane, cv::Rect(0, 3 * fh, fw, fh)); - - cv::gpu::cartToPolar(dfdx, dfdy, mag, ang, true, s); - - // normolize magnitude to uchar interval and angles to 6 bins - GpuMat nmag(fplane, cv::Rect(0, 4 * fh, fw, fh)); - GpuMat nang(fplane, cv::Rect(0, 5 * fh, fw, fh)); - - cv::gpu::multiply(mag, cv::Scalar::all(1.f / (8 *::log(2.0f))), nmag, 1, -1, s); - cv::gpu::multiply(ang, cv::Scalar::all(1.f / 60.f), nang, 1, -1, s); - - //create uchar magnitude - GpuMat cmag(plane, cv::Rect(0, fh * Fields::HOG_BINS, fw, fh)); - if (s) - s.enqueueConvert(nmag, cmag, CV_8UC1); - else - nmag.convertTo(cmag, CV_8UC1); - - cudaStream_t stream = StreamAccessor::getStream(s); - device::icf::fillBins(plane, nang, fw, fh, Fields::HOG_BINS, stream); - } - - void createLuvBins(const cv::gpu::GpuMat& colored, Stream& s) - { - static const int fw = colored.cols; - static const int fh = colored.rows; - - cv::gpu::cvtColor(colored, luv, CV_BGR2Luv, s); - - std::vector splited; - for(int i = 0; i < Fields::LUV_BINS; ++i) - { - splited.push_back(GpuMat(plane, cv::Rect(0, fh * (7 + i), fw, fh))); - } - - cv::gpu::split(luv, splited, s); - } - - void integrate(const int fh, const int fw, Stream& s) - { - GpuMat channels(plane, cv::Rect(0, 0, fw, fh * Fields::HOG_LUV_BINS)); - cv::gpu::resize(channels, shrunk, cv::Size(), 1.f / shrinkage, 1.f / shrinkage, CV_INTER_AREA, s); - - if (info.majorVersion() < 3) - cv::gpu::integralBuffered(shrunk, hogluv, integralBuffer, s); - else - { - cudaStream_t stream = StreamAccessor::getStream(s); - device::imgproc::shfl_integral_gpu_buffered(shrunk, integralBuffer, hogluv, 12, stream); - } - } - public: + cv::Ptr preprocessor; + // scales range float minScale; float maxScale; @@ -494,14 +404,6 @@ public: const int shrinkage; int downscales; - // preallocated buffer 640x480x10 for hogluv + 640x480 got gray - GpuMat plane; - - // preallocated buffer for floating point operations - GpuMat fplane; - - // temporial mat for cvtColor - GpuMat luv; // 160x120x10 GpuMat shrunk; @@ -512,11 +414,12 @@ public: // 161x121x10 GpuMat hogluv; - // used for area overlap computing during - GpuMat overlaps; // used for suppression GpuMat suppressed; + // used for area overlap computing during + GpuMat overlaps; + // Cascade from xml GpuMat octaves; @@ -525,36 +428,36 @@ public: GpuMat leaves; GpuMat levels; - GpuMat sobelBuf; - GpuMat collected; + // For ROI + GpuMat mask; + GpuMat genRoiTmp; + +// GpuMat collected; - cv::gpu::GpuMat genRoiTmp; std::vector voctaves; - DeviceInfo info; +// DeviceInfo info; enum { BOOST = 0 }; enum { - FRAME_WIDTH = 640, - FRAME_HEIGHT = 480, - HOG_BINS = 6, - LUV_BINS = 3, - HOG_LUV_BINS = 10 + DEFAULT_FRAME_WIDTH = 640, + DEFAULT_FRAME_HEIGHT = 480, + HOG_LUV_BINS = 10 }; }; -cv::gpu::SCascade::SCascade(const double mins, const double maxs, const int sc, const int rjf) -: fields(0), minScale(mins), maxScale(maxs), scales(sc), rejCriteria(rjf) {} +cv::gpu::SCascade::SCascade(const double mins, const double maxs, const int sc, const int fl) +: fields(0), minScale(mins), maxScale(maxs), scales(sc), flags(fl) {} cv::gpu::SCascade::~SCascade() { delete fields; } bool cv::gpu::SCascade::load(const FileNode& fn) { if (fields) delete fields; - fields = Fields::parseCascade(fn, minScale, maxScale, scales); + fields = Fields::parseCascade(fn, minScale, maxScale, scales, flags); return fields != 0; } @@ -572,12 +475,24 @@ void cv::gpu::SCascade::detect(InputArray _image, InputArray _rois, OutputArray GpuMat rois = _rois.getGpuMat(), objects = _objects.getGpuMat(); + /// roi Fields& flds = *fields; + int shr = flds.shrinkage; + + flds.mask.create( rois.cols / shr, rois.rows / shr, rois.type()); + + cv::gpu::resize(rois, flds.genRoiTmp, cv::Size(), 1.f / shr, 1.f / shr, CV_INTER_AREA, s); + cv::gpu::transpose(flds.genRoiTmp, flds.mask, s); + if (type == CV_8UC3) { - if (!flds.update(image.rows, image.cols, flds.shrinkage) || flds.check(minScale, maxScale, scales)) + flds.update(image.rows, image.cols, flds.shrinkage); + + if (flds.check(minScale, maxScale, scales)) flds.createLevels(image.rows, image.cols); - flds.preprocess(image, s); + + flds.preprocessor->apply(image, flds.shrunk); + cv::gpu::integralBuffered(flds.shrunk, flds.hogluv, flds.integralBuffer, s); } else { @@ -587,9 +502,9 @@ void cv::gpu::SCascade::detect(InputArray _image, InputArray _rois, OutputArray image.copyTo(flds.hogluv); } - flds.detect(rois, objects, s); + flds.detect(objects, s); - if (rejCriteria != NO_REJECT) + if ( (flags && NMS_MASK) != NO_REJECT) { GpuMat spr(objects, cv::Rect(0, 0, flds.suppressed.cols, flds.suppressed.rows)); flds.suppress(objects, s); @@ -597,79 +512,122 @@ void cv::gpu::SCascade::detect(InputArray _image, InputArray _rois, OutputArray } } -void cv::gpu::SCascade::genRoi(InputArray _roi, OutputArray _mask, Stream& stream) const +void cv::gpu::SCascade::read(const FileNode& fn) { - CV_Assert(fields); - int shr = (*fields).shrinkage; + Algorithm::read(fn); +} - const GpuMat roi = _roi.getGpuMat(); - _mask.create( roi.cols / shr, roi.rows / shr, roi.type()); - GpuMat mask = _mask.getGpuMat(); +namespace { - GpuMat& tmp = (*fields).genRoiTmp; - cv::gpu::resize(roi, tmp, cv::Size(), 1.f / shr, 1.f / shr, CV_INTER_AREA, stream); - cv::gpu::transpose(tmp, mask, stream); -} +using cv::InputArray; +using cv::OutputArray; +using cv::gpu::Stream; +using cv::gpu::GpuMat; -void cv::gpu::SCascade::read(const FileNode& fn) +inline void setZero(cv::gpu::GpuMat& m, Stream& s) { - Algorithm::read(fn); + if (s) + s.enqueueMemSet(m, 0); + else + m.setTo(0); } -// namespace { +struct GenricPreprocessor : public cv::gpu::ChannelsProcessor +{ + GenricPreprocessor(const int s, const int b) : cv::gpu::ChannelsProcessor(), shrinkage(s), bins(b) {} + virtual ~GenricPreprocessor() {} -// void bgr2Luv(const cv::gpu::GpuMat& input, cv::gpu::GpuMat& luv /*integral*/) -// { -// cv::gpu::GpuMat bgr; -// cv::gpu::GaussianBlur(input, bgr, cv::Size(3, 3), -1); + virtual void apply(InputArray _frame, OutputArray _shrunk, Stream& s = Stream::Null()) + { + const GpuMat frame = _frame.getGpuMat(); -// cv::gpu::GpuMat gray, /*luv,*/ shrunk, buffer; -// luv.create(bgr.rows * 10, bgr.cols, CV_8UC1); -// luv.setTo(0); + _shrunk.create(frame.rows * (4 + bins) / shrinkage, frame.cols / shrinkage, CV_8UC1); + GpuMat shrunk = _shrunk.getGpuMat(); -// cv::gpu::cvtColor(bgr, gray, CV_BGR2GRAY); -// cv::gpu::device::icf::magnitude(gray, luv(cv::Rect(0, 0, bgr.cols, bgr.rows * 7))); + channels.create(frame.rows * (4 + bins), frame.cols, CV_8UC1); + setZero(channels, s); -// cv::gpu::GpuMat __luv(luv, cv::Rect(0, bgr.rows * 7, bgr.cols, bgr.rows * 3)); -// cv::gpu::device::icf::bgr2Luv(bgr, __luv); + cv::gpu::cvtColor(frame, gray, CV_BGR2GRAY, s); + createHogBins(s); -// // cv::gpu::resize(luv, shrunk, cv::Size(), 0.25f, 0.25f, CV_INTER_AREA); -// // cv::gpu::integralBuffered(shrunk, integral, buffer); -// } -// } + createLuvBins(frame, s); -namespace { + cv::gpu::resize(channels, shrunk, cv::Size(), 1.f / shrinkage, 1.f / shrinkage, CV_INTER_AREA, s); + } -using cv::InputArray; -using cv::OutputArray; -using cv::gpu::Stream; -using cv::gpu::GpuMat; +private: -struct GenricPreprocessor : public cv::gpu::SCascade::Preprocessor -{ - GenricPreprocessor(const int s, const int b) : cv::gpu::SCascade::Preprocessor(), shrinkage(s), bins(b) {} + void createHogBins(Stream& s) + { + static const int fw = gray.cols; + static const int fh = gray.rows; + + fplane.create(fh * HOG_BINS, fw, CV_32FC1); + + GpuMat dfdx(fplane, cv::Rect(0, 0, fw, fh)); + GpuMat dfdy(fplane, cv::Rect(0, fh, fw, fh)); - virtual void apply(InputArray /*frame*/, OutputArray /*channels*/, Stream& /*s*/ = Stream::Null()) + cv::gpu::Sobel(gray, dfdx, CV_32F, 1, 0, sobelBuf, 3, 1, cv::BORDER_DEFAULT, -1, s); + cv::gpu::Sobel(gray, dfdy, CV_32F, 0, 1, sobelBuf, 3, 1, cv::BORDER_DEFAULT, -1, s); + + GpuMat mag(fplane, cv::Rect(0, 2 * fh, fw, fh)); + GpuMat ang(fplane, cv::Rect(0, 3 * fh, fw, fh)); + + cv::gpu::cartToPolar(dfdx, dfdy, mag, ang, true, s); + + // normolize magnitude to uchar interval and angles to 6 bins + GpuMat nmag(fplane, cv::Rect(0, 4 * fh, fw, fh)); + GpuMat nang(fplane, cv::Rect(0, 5 * fh, fw, fh)); + + cv::gpu::multiply(mag, cv::Scalar::all(1.f / (8 *::log(2.0f))), nmag, 1, -1, s); + cv::gpu::multiply(ang, cv::Scalar::all(1.f / 60.f), nang, 1, -1, s); + + //create uchar magnitude + GpuMat cmag(channels, cv::Rect(0, fh * HOG_BINS, fw, fh)); + if (s) + s.enqueueConvert(nmag, cmag, CV_8UC1); + else + nmag.convertTo(cmag, CV_8UC1); + + cudaStream_t stream = cv::gpu::StreamAccessor::getStream(s); + cv::gpu::device::icf::fillBins(channels, nang, fw, fh, HOG_BINS, stream); + } + + void createLuvBins(const cv::gpu::GpuMat& colored, Stream& s) { + static const int fw = colored.cols; + static const int fh = colored.rows; + + cv::gpu::cvtColor(colored, luv, CV_BGR2Luv, s); + std::vector splited; + for(int i = 0; i < LUV_BINS; ++i) + { + splited.push_back(GpuMat(channels, cv::Rect(0, fh * (7 + i), fw, fh))); + } + + cv::gpu::split(luv, splited, s); } -private: + enum {HOG_BINS = 6, LUV_BINS = 3}; + const int shrinkage; const int bins; + + GpuMat gray; + GpuMat luv; + GpuMat channels; + + // preallocated buffer for floating point operations + GpuMat fplane; + GpuMat sobelBuf; }; -inline void setZero(cv::gpu::GpuMat& m, Stream& s) -{ - if (s) - s.enqueueMemSet(m, 0); - else - m.setTo(0); -} -struct SeparablePreprocessor : public cv::gpu::SCascade::Preprocessor +struct SeparablePreprocessor : public cv::gpu::ChannelsProcessor { - SeparablePreprocessor(const int s, const int b) : cv::gpu::SCascade::Preprocessor(), shrinkage(s), bins(b) {} + SeparablePreprocessor(const int s, const int b) : cv::gpu::ChannelsProcessor(), shrinkage(s), bins(b) {} + virtual ~SeparablePreprocessor() {} virtual void apply(InputArray _frame, OutputArray _shrunk, Stream& s = Stream::Null()) { @@ -701,16 +659,17 @@ private: } -cv::gpu::SCascade::Preprocessor::Preprocessor(){} - -cv::Ptr cv::gpu::SCascade::Preprocessor::create(const int s, const int b, const int m) +cv::Ptr cv::gpu::ChannelsProcessor::create(const int s, const int b, const int m) { - CV_Assert(m == SEPARABLE || m == GENERIC); + CV_Assert((m && SEPARABLE) || (m && GENERIC)); - if (m == GENERIC) - return cv::Ptr(new GenricPreprocessor(s, b)); + if (m && GENERIC) + return cv::Ptr(new GenricPreprocessor(s, b)); - return cv::Ptr(new SeparablePreprocessor(s, b)); + return cv::Ptr(new SeparablePreprocessor(s, b)); } +cv::gpu::ChannelsProcessor::ChannelsProcessor() { } +cv::gpu::ChannelsProcessor::~ChannelsProcessor() { } + #endif \ No newline at end of file diff --git a/modules/gpu/test/test_softcascade.cpp b/modules/gpu/test/test_softcascade.cpp index 266a7d4..28ef15b 100644 --- a/modules/gpu/test/test_softcascade.cpp +++ b/modules/gpu/test/test_softcascade.cpp @@ -169,7 +169,7 @@ GPU_TEST_P(SCascadeTestRoi, detect, ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); - GpuMat colored(coloredCpu), objectBoxes(1, 16384, CV_8UC1), rois(colored.size(), CV_8UC1), trois; + GpuMat colored(coloredCpu), objectBoxes(1, 16384, CV_8UC1), rois(colored.size(), CV_8UC1); rois.setTo(0); int nroi = GET_PARAM(3); @@ -183,8 +183,8 @@ GPU_TEST_P(SCascadeTestRoi, detect, cv::rectangle(result, r, cv::Scalar(0, 0, 255, 255), 1); } objectBoxes.setTo(0); - cascade.genRoi(rois, trois); - cascade.detect(colored, trois, objectBoxes); + + cascade.detect(colored, rois, objectBoxes); cv::Mat dt(objectBoxes); typedef cv::gpu::SCascade::Detection Detection; @@ -239,10 +239,8 @@ GPU_TEST_P(SCascadeTestAll, detect, GpuMat sub(rois, cv::Rect(rois.cols / 4, rois.rows / 4,rois.cols / 2, rois.rows / 2)); sub.setTo(cv::Scalar::all(1)); - cv::gpu::GpuMat trois; - cascade.genRoi(rois, trois); objectBoxes.setTo(0); - cascade.detect(colored, trois, objectBoxes); + cascade.detect(colored, rois, objectBoxes); typedef cv::gpu::SCascade::Detection Detection; cv::Mat detections(objectBoxes); @@ -279,10 +277,8 @@ GPU_TEST_P(SCascadeTestAll, detectOnIntegral, GpuMat objectBoxes(1, 100000, CV_8UC1), rois(cv::Size(640, 480), CV_8UC1); rois.setTo(1); - cv::gpu::GpuMat trois; - cascade.genRoi(rois, trois); objectBoxes.setTo(0); - cascade.detect(hogluv, trois, objectBoxes); + cascade.detect(hogluv, rois, objectBoxes); typedef cv::gpu::SCascade::Detection Detection; cv::Mat detections(objectBoxes); @@ -315,12 +311,9 @@ GPU_TEST_P(SCascadeTestAll, detectStream, cv::gpu::Stream s; - cv::gpu::GpuMat trois; - cascade.genRoi(rois, trois, s); objectBoxes.setTo(0); - cascade.detect(colored, trois, objectBoxes, s); - - cudaDeviceSynchronize(); + cascade.detect(colored, rois, objectBoxes, s); + s.waitForCompletion(); typedef cv::gpu::SCascade::Detection Detection; cv::Mat detections(objectBoxes); -- 2.7.4