From 267d140bfeb43b1c47734e554c37c0b1db81787f Mon Sep 17 00:00:00 2001
From: "marina.kolpakova" <marina.kolpakova@itseez.com>
Date: Thu, 20 Sep 2012 16:22:10 +0400
Subject: [PATCH] soft cascade: gpu representation
---
modules/gpu/include/opencv2/gpu/gpu.hpp | 9 +-
modules/gpu/src/cuda/isf-sc.cu | 43 +++++
modules/gpu/src/icf.hpp | 118 ++++++++++++
modules/gpu/src/softcascade.cpp | 236 +++++++++++++++++++++++-
modules/gpu/test/test_softcascade.cpp | 73 ++++++++
5 files changed, 473 insertions(+), 6 deletions(-)
create mode 100644 modules/gpu/src/cuda/isf-sc.cu
create mode 100644 modules/gpu/src/icf.hpp
create mode 100644 modules/gpu/test/test_softcascade.cpp
diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 4a2d88aa0..61f6006c5 100644
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -1554,9 +1554,14 @@ public:
virtual ~SoftCascade();
- //! return vector of bounding boxes. Each box contains one detected object
+ //! detect specific objects on in the input frame for all scales computed flom minScale and maxscale values
+ //! Param image is input frame for detector. Cascade will be applied to it.
+ //! Param rois is a mask
+ //! Param objects 4-channel matrix thet contain detected rectangles
+ //! Param rejectfactor used for final object box computing
+ //! Param stream
virtual void detectMultiScale(const GpuMat& image, const GpuMat& rois, GpuMat& objects,
- int rejectfactor = 1, Stream stream = Stream::Null()); // ToDo store objects in GPU mem
+ int rejectfactor = 1, Stream stream = Stream::Null());
protected:
enum { BOOST = 0 };
diff --git a/modules/gpu/src/cuda/isf-sc.cu b/modules/gpu/src/cuda/isf-sc.cu
new file mode 100644
index 000000000..f36f86f96
--- /dev/null
+++ b/modules/gpu/src/cuda/isf-sc.cu
@@ -0,0 +1,43 @@
+/*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.
+//
+//
+// License Agreement
+// For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2008-2012, Willow Garage Inc., 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 the copyright holders 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 <icf.hpp>
\ No newline at end of file
diff --git a/modules/gpu/src/icf.hpp b/modules/gpu/src/icf.hpp
new file mode 100644
index 000000000..110890232
--- /dev/null
+++ b/modules/gpu/src/icf.hpp
@@ -0,0 +1,118 @@
+/*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.
+//
+//
+// License Agreement
+// For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2008-2012, Willow Garage Inc., 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 the copyright holders 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*/
+
+#ifndef __OPENCV_ICF_HPP__
+#define __OPENCV_ICF_HPP__
+
+#if defined __CUDACC__
+# define __hd__ __host__ __device__ __forceinline__
+#else
+# define __hd__
+#endif
+
+
+namespace icf {
+
+ struct Cascade
+ {
+
+ };
+
+ struct ChannelStorage
+ {
+
+ };
+
+ struct __align__(16) Octave
+ {
+ ushort index;
+ ushort stages;
+ ushort shrinkage;
+ ushort2 size;
+ float scale;
+
+ Octave(const ushort i, const ushort s, const ushort sh, const ushort2 sz, const float sc)
+ : index(i), stages(s), shrinkage(sh), size(sz), scale(sc) {}
+ };
+
+ struct __align__(8) Node
+ {
+ int feature;
+ float threshold;
+
+ Node(const int f, const float t) : feature(f), threshold(t) {}
+ };
+
+ struct __align__(8) Feature
+ {
+ int channel;
+ uchar4 rect;
+
+ Feature(const int c, const uchar4 r) : channel(c), rect(r) {}
+ };
+
+ struct __align__(8) Level //is actually 24 bytes
+ {
+ int octave;
+
+ // float origScale; //not actually used
+ float relScale;
+ float shrScale; // used for marking detection
+ float scaling[2]; // calculated according to Dollal paper
+
+ // for 640x480 we can not get overflow
+ uchar2 workRect;
+ uchar2 objSize;
+
+ Level(int idx, const Octave& oct, const float scale, const int w, const int h)
+ : octave(idx), relScale(scale / oct.scale), shrScale (relScale / (float)oct.shrinkage)
+ {
+ workRect.x = round(w / (float)oct.shrinkage);
+ workRect.y = round(h / (float)oct.shrinkage);
+
+ objSize.x = round(oct.size.x * relScale);
+ objSize.y = round(oct.size.y * relScale);
+ }
+ };
+}
+
+#endif
\ No newline at end of file
diff --git a/modules/gpu/src/softcascade.cpp b/modules/gpu/src/softcascade.cpp
index 509e3f501..04b68539c 100644
--- a/modules/gpu/src/softcascade.cpp
+++ b/modules/gpu/src/softcascade.cpp
@@ -56,12 +56,242 @@ void cv::gpu::SoftCascade::detectMultiScale(const GpuMat&, const GpuMat&, GpuMat
#else
+#include <icf.hpp>
+
struct cv::gpu::SoftCascade::Filds
{
- bool fill(const FileNode &root, const float mins, const float maxs){return true;}
- void calcLevels(int frameW, int frameH, int scales) {}
+ // scales range
+ float minScale;
+ float maxScale;
+
+ int origObjWidth;
+ int origObjHeight;
+
+ GpuMat octaves;
+ GpuMat stages;
+ GpuMat nodes;
+ GpuMat leaves;
+ GpuMat features;
+
+ std::vector<float> scales;
+
+ icf::Cascade cascade;
+
+ bool fill(const FileNode &root, const float mins, const float maxs);
+
+private:
+ void calcLevels(const std::vector<icf::Octave>& octs,
+ int frameW, int frameH, int nscales);
+
+ typedef std::vector<icf::Octave>::const_iterator octIt_t;
+ int fitOctave(const std::vector<icf::Octave>& octs, const float& logFactor)
+ {
+ float minAbsLog = FLT_MAX;
+ int res = 0;
+ for (int oct = 0; oct < (int)octs.size(); ++oct)
+ {
+ const icf::Octave& octave =octs[oct];
+ float logOctave = ::log(octave.scale);
+ float logAbsScale = ::fabs(logFactor - logOctave);
+
+ if(logAbsScale < minAbsLog)
+ {
+ res = oct;
+ minAbsLog = logAbsScale;
+ }
+ }
+ return res;
+ }
};
+inline bool cv::gpu::SoftCascade::Filds::fill(const FileNode &root, const float mins, const float maxs)
+{
+ minScale = mins;
+ maxScale = maxs;
+
+ // cascade properties
+ static const char *const SC_STAGE_TYPE = "stageType";
+ static const char *const SC_BOOST = "BOOST";
+
+ static const char *const SC_FEATURE_TYPE = "featureType";
+ static const char *const SC_ICF = "ICF";
+
+ static const char *const SC_ORIG_W = "width";
+ static const char *const SC_ORIG_H = "height";
+
+ static const char *const SC_OCTAVES = "octaves";
+ static const char *const SC_STAGES = "stages";
+ static const char *const SC_FEATURES = "features";
+
+ static const char *const SC_WEEK = "weakClassifiers";
+ static const char *const SC_INTERNAL = "internalNodes";
+ static const char *const SC_LEAF = "leafValues";
+
+ static const char *const SC_OCT_SCALE = "scale";
+ static const char *const SC_OCT_STAGES = "stageNum";
+ static const char *const SC_OCT_SHRINKAGE = "shrinkingFactor";
+
+ static const char *const SC_STAGE_THRESHOLD = "stageThreshold";
+
+ static const char * const SC_F_CHANNEL = "channel";
+ static const char * const SC_F_RECT = "rect";
+
+ // only Ada Boost supported
+ std::string stageTypeStr = (string)root[SC_STAGE_TYPE];
+ CV_Assert(stageTypeStr == SC_BOOST);
+
+ // only HOG-like integral channel features cupported
+ string featureTypeStr = (string)root[SC_FEATURE_TYPE];
+ CV_Assert(featureTypeStr == SC_ICF);
+
+ origObjWidth = (int)root[SC_ORIG_W];
+ CV_Assert(origObjWidth == SoftCascade::ORIG_OBJECT_WIDTH);
+
+ origObjHeight = (int)root[SC_ORIG_H];
+ CV_Assert(origObjHeight == SoftCascade::ORIG_OBJECT_HEIGHT);
+
+ FileNode fn = root[SC_OCTAVES];
+ if (fn.empty()) return false;
+
+ std::vector<icf::Octave> voctaves;
+ std::vector<float> vstages;
+ std::vector<icf::Node> vnodes;
+ std::vector<float> vleaves;
+ std::vector<icf::Feature> vfeatures;
+ scales.clear();
+
+ // std::vector<Level> levels;
+
+ FileNodeIterator it = fn.begin(), it_end = fn.end();
+ int feature_offset = 0;
+ ushort octIndex = 0;
+
+ for (; it != it_end; ++it)
+ {
+ FileNode fns = *it;
+ float scale = (float)fns[SC_OCT_SCALE];
+ scales.push_back(scale);
+ ushort nstages = saturate_cast<ushort>((int)fn[SC_OCT_STAGES]);
+ ushort2 size;
+ size.x = cvRound(SoftCascade::ORIG_OBJECT_WIDTH * scale);
+ size.y = cvRound(SoftCascade::ORIG_OBJECT_HEIGHT * scale);
+ ushort shrinkage = saturate_cast<ushort>((int)fn[SC_OCT_SHRINKAGE]);
+
+ icf::Octave octave(octIndex, nstages, shrinkage, size, scale);
+ CV_Assert(octave.stages > 0);
+ voctaves.push_back(octave);
+
+ FileNode ffs = fns[SC_FEATURES];
+ if (ffs.empty()) return false;
+
+ fns = fns[SC_STAGES];
+ if (fn.empty()) return false;
+
+ // for each stage (~ decision tree with H = 2)
+ FileNodeIterator st = fns.begin(), st_end = fns.end();
+ for (; st != st_end; ++st )
+ {
+ fns = *st;
+ vstages.push_back((float)fn[SC_STAGE_THRESHOLD]);
+
+ fns = fns[SC_WEEK];
+ FileNodeIterator ftr = fns.begin(), ft_end = fns.end();
+ for (; ftr != ft_end; ++ftr)
+ {
+ fns = (*ftr)[SC_INTERNAL];
+ FileNodeIterator inIt = fns.begin(), inIt_end = fns.end();
+ for (; inIt != inIt_end;)
+ {
+ int feature = (int)(*(inIt +=2)++) + feature_offset;
+ vnodes.push_back(icf::Node(feature, (float)(*(inIt++))));
+ }
+
+ fns = (*ftr)[SC_LEAF];
+ inIt = fns.begin(), inIt_end = fns.end();
+ for (; inIt != inIt_end; ++inIt)
+ vleaves.push_back((float)(*inIt));
+ }
+ }
+
+ st = ffs.begin(), st_end = ffs.end();
+ for (; st != st_end; ++st )
+ {
+ cv::FileNode rn = (*st)[SC_F_RECT];
+ cv::FileNodeIterator r_it = rn.begin();
+ uchar4 rect;
+ rect.x = saturate_cast<uchar>((int)*(r_it++));
+ rect.y = saturate_cast<uchar>((int)*(r_it++));
+ rect.z = saturate_cast<uchar>((int)*(r_it++));
+ rect.w = saturate_cast<uchar>((int)*(r_it++));
+ vfeatures.push_back(icf::Feature((int)(*st)[SC_F_CHANNEL], rect));
+ }
+
+ feature_offset += octave.stages * 3;
+ ++octIndex;
+ }
+
+ // upload in gpu memory
+ octaves.upload(cv::Mat(1, voctaves.size() * sizeof(icf::Octave), CV_8UC1, (uchar*)&(voctaves[0]) ));
+ CV_Assert(!octaves.empty());
+
+ stages.upload(cv::Mat(vstages).reshape(1,1));
+ CV_Assert(!stages.empty());
+
+ nodes.upload(cv::Mat(1, vnodes.size() * sizeof(icf::Node), CV_8UC1, (uchar*)&(vnodes[0]) ));
+ CV_Assert(!nodes.empty());
+
+ leaves.upload(cv::Mat(vleaves).reshape(1,1));
+ CV_Assert(!leaves.empty());
+
+ features.upload(cv::Mat(1, vfeatures.size() * sizeof(icf::Feature), CV_8UC1, (uchar*)&(vfeatures[0]) ));
+ CV_Assert(!features.empty());
+
+ // compute levels
+ calcLevels(voctaves, (int)SoftCascade::FRAME_WIDTH, (int)SoftCascade::FRAME_HEIGHT, (int)SoftCascade::TOTAL_SCALES);
+
+ return true;
+}
+
+inline void cv::gpu::SoftCascade::Filds::calcLevels(const std::vector<icf::Octave>& octs,
+ int frameW, int frameH, int nscales)
+{
+ CV_Assert(nscales > 1);
+
+ std::vector<icf::Level> levels;
+ float logFactor = (::log(maxScale) - ::log(minScale)) / (nscales -1);
+
+ float scale = minScale;
+ for (int sc = 0; sc < nscales; ++sc)
+ {
+ int width = ::std::max(0.0f, frameW - (origObjWidth * scale));
+ int height = ::std::max(0.0f, frameH - (origObjHeight * scale));
+
+ float logScale = ::log(scale);
+ int fit = fitOctave(octs, logScale);
+
+ icf::Level level(fit, octs[fit], scale, width, height);
+
+ if (!width || !height)
+ break;
+ else
+ levels.push_back(level);
+
+ if (::fabs(scale - maxScale) < FLT_EPSILON) break;
+ scale = ::std::min(maxScale, ::expf(::log(scale) + logFactor));
+
+ // std::cout << "level " << sc << " scale "
+ // << levels[sc].origScale
+ // << " octeve "
+ // << levels[sc].octave->scale
+ // << " "
+ // << levels[sc].relScale
+ // << " " << levels[sc].shrScale
+ // << " [" << levels[sc].objSize.width
+ // << " " << levels[sc].objSize.height << "] ["
+ // << levels[sc].workRect.width << " " << levels[sc].workRect.height << "]" << std::endl;
+ }
+}
+
cv::gpu::SoftCascade::SoftCascade() : filds(0) {}
cv::gpu::SoftCascade::SoftCascade( const string& filename, const float minScale, const float maxScale) : filds(0)
@@ -86,8 +316,6 @@ bool cv::gpu::SoftCascade::load( const string& filename, const float minScale, c
filds = new Filds;
Filds& flds = *filds;
if (!flds.fill(fs.getFirstTopLevelNode(), minScale, maxScale)) return false;
- flds.calcLevels(FRAME_WIDTH, FRAME_HEIGHT, TOTAL_SCALES);
-
return true;
}
diff --git a/modules/gpu/test/test_softcascade.cpp b/modules/gpu/test/test_softcascade.cpp
new file mode 100644
index 000000000..821a2b140
--- /dev/null
+++ b/modules/gpu/test/test_softcascade.cpp
@@ -0,0 +1,73 @@
+/*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.
+//
+//
+// License Agreement
+// For Open Source Computer Vision Library
+//
+// Copyright (C) 2000-2008, Intel Corporation, all rights reserved.
+// Copyright (C) 2008-2012, Willow Garage Inc., 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 the copyright holders 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 <test_precomp.hpp>
+
+#ifdef HAVE_CUDA
+
+using cv::gpu::GpuMat;
+
+TEST(SoftCascade, readCascade)
+{
+ std::string xml = cvtest::TS::ptr()->get_data_path() + "cascadeandhog/icf-template.xml";
+ cv::gpu::SoftCascade cascade;
+ ASSERT_TRUE(cascade.load(xml));
+
+}
+
+TEST(SoftCascade, detect)
+{
+ std::string xml = cvtest::TS::ptr()->get_data_path() + "cascadeandhog/sc_cvpr_2012_to_opencv.xml";
+ cv::gpu::SoftCascade cascade;
+ ASSERT_TRUE(cascade.load(xml));
+
+ cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() + "cascadeandhog/bahnhof/image_00000000_0.png");
+ ASSERT_FALSE(coloredCpu.empty());
+ GpuMat colored(coloredCpu), objectBoxes, rois;
+
+ // ASSERT_NO_THROW(
+ // {
+ cascade.detectMultiScale(colored, rois, objectBoxes);
+ // });
+}
+
+#endif
\ No newline at end of file