generic-library/opencv
1
0
Fork 0
Code Issues Pull Requests Releases Wiki Activity

integrate new cascade format to GPU soft cascade implementation

Browse Source
This commit is contained in:
marina.kolpakova 2013-01-23 14:57:51 +04:00
parent 7572b4d400
commit 8d9c9c2690
7 changed files with 210 additions and 156 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

8
modules/gpu/include/opencv2/gpu/gpu.hpp
View File

@ -1556,7 +1556,7 @@ protected:
ChannelsProcessor(); ChannelsProcessor();
}; };
// Implementation of soft (stageless) cascaded detector. // Implementation of soft (stage-less) cascaded detector.
class CV_EXPORTS SCascade : public Algorithm class CV_EXPORTS SCascade : public Algorithm
{ {
public: public:
@ -1577,8 +1577,8 @@ public:
enum { NO_REJECT = 1, DOLLAR = 2, /*PASCAL = 4,*/ DEFAULT = NO_REJECT, NMS_MASK = 0xF}; enum { NO_REJECT = 1, DOLLAR = 2, /*PASCAL = 4,*/ DEFAULT = NO_REJECT, NMS_MASK = 0xF};
// An empty cascade will be created. // 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 minimum scale relative to the original size of the image on which cascade will be applied.
// Param minScale is a maximum 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 applied.
// Param scales is a number of scales from minScale to maxScale. // Param scales is a number of scales from minScale to maxScale.
// Param flags is an extra tuning flags. // Param flags is an extra tuning flags.
SCascade(const double minScale = 0.4, const double maxScale = 5., const int scales = 55, SCascade(const double minScale = 0.4, const double maxScale = 5., const int scales = 55,
@ -1595,7 +1595,7 @@ public:
// Load cascade config. // Load cascade config.
virtual void read(const FileNode& fn); virtual void read(const FileNode& fn);
// Return the matrix of of detectioned objects. // Return the matrix of of detected objects.
// Param image is a frame on which detector will be applied. // Param image is a frame on which detector will be applied.
// Param rois is a regions of interests mask generated by genRoi. // Param rois is a regions of interests mask generated by genRoi.
// Only the objects that fall into one of the regions will be returned. // Only the objects that fall into one of the regions will be returned.

56
modules/gpu/perf/perf_softcascade.cpp
View File

@ -1,6 +1,6 @@
#include "perf_precomp.hpp" #include "perf_precomp.hpp"
#define PERF_TEST_P1(fixture, name, params) \ #define SC_PERF_TEST_P(fixture, name, params) \
class fixture##_##name : public fixture {\ class fixture##_##name : public fixture {\
public:\ public:\
fixture##_##name() {}\ fixture##_##name() {}\
@ -28,7 +28,7 @@ namespace {
bool operator()(const cv::gpu::SCascade::Detection& a, bool operator()(const cv::gpu::SCascade::Detection& a,
const cv::gpu::SCascade::Detection& b) const const cv::gpu::SCascade::Detection& b) const
{ {
if (a.x != b.x) return a.x < b.x; if (a.x != b.x) return a.x < b.x;
else if (a.y != b.y) return a.y < b.y; else if (a.y != b.y) return a.y < b.y;
else if (a.w != b.w) return a.w < b.w; else if (a.w != b.w) return a.w < b.w;
else return a.h < b.h; else return a.h < b.h;
@ -52,10 +52,11 @@ namespace {
typedef std::tr1::tuple<std::string, std::string> fixture_t; typedef std::tr1::tuple<std::string, std::string> fixture_t;
typedef perf::TestBaseWithParam<fixture_t> SCascadeTest; typedef perf::TestBaseWithParam<fixture_t> SCascadeTest;
PERF_TEST_P1(SCascadeTest, detect, SC_PERF_TEST_P(SCascadeTest, detect,
testing::Combine( testing::Combine(
testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), testing::Values(std::string("cv/cascadeandhog/cascades/inria_caltech-17.01.2013.xml"),
testing::Values(std::string("cv/cascadeandhog/bahnhof/image_00000000_0.png")))) std::string("cv/cascadeandhog/cascades/sc_cvpr_2012_to_opencv_new_format.xml")),
testing::Values(std::string("cv/cascadeandhog/images/image_00000000_0.png"))))
RUN_GPU(SCascadeTest, detect) RUN_GPU(SCascadeTest, detect)
{ {
@ -108,10 +109,11 @@ static cv::Rect getFromTable(int idx)
typedef std::tr1::tuple<std::string, std::string, int> roi_fixture_t; typedef std::tr1::tuple<std::string, std::string, int> roi_fixture_t;
typedef perf::TestBaseWithParam<roi_fixture_t> SCascadeTestRoi; typedef perf::TestBaseWithParam<roi_fixture_t> SCascadeTestRoi;
PERF_TEST_P1(SCascadeTestRoi, detectInRoi, SC_PERF_TEST_P(SCascadeTestRoi, detectInRoi,
testing::Combine( testing::Combine(
testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), testing::Values(std::string("cv/cascadeandhog/cascades/inria_caltech-17.01.2013.xml"),
testing::Values(std::string("cv/cascadeandhog/bahnhof/image_00000000_0.png")), std::string("cv/cascadeandhog/cascades/sc_cvpr_2012_to_opencv_new_format.xml")),
testing::Values(std::string("cv/cascadeandhog/images/image_00000000_0.png")),
testing::Range(0, 5))) testing::Range(0, 5)))
RUN_GPU(SCascadeTestRoi, detectInRoi) RUN_GPU(SCascadeTestRoi, detectInRoi)
@ -152,10 +154,11 @@ RUN_GPU(SCascadeTestRoi, detectInRoi)
NO_CPU(SCascadeTestRoi, detectInRoi) NO_CPU(SCascadeTestRoi, detectInRoi)
PERF_TEST_P1(SCascadeTestRoi, detectEachRoi, SC_PERF_TEST_P(SCascadeTestRoi, detectEachRoi,
testing::Combine( testing::Combine(
testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), testing::Values(std::string("cv/cascadeandhog/cascades/inria_caltech-17.01.2013.xml"),
testing::Values(std::string("cv/cascadeandhog/bahnhof/image_00000000_0.png")), std::string("cv/cascadeandhog/cascades/sc_cvpr_2012_to_opencv_new_format.xml")),
testing::Values(std::string("cv/cascadeandhog/images/image_00000000_0.png")),
testing::Range(0, 10))) testing::Range(0, 10)))
RUN_GPU(SCascadeTestRoi, detectEachRoi) RUN_GPU(SCascadeTestRoi, detectEachRoi)
@ -191,9 +194,10 @@ RUN_GPU(SCascadeTestRoi, detectEachRoi)
NO_CPU(SCascadeTestRoi, detectEachRoi) NO_CPU(SCascadeTestRoi, detectEachRoi)
PERF_TEST_P1(SCascadeTest, detectOnIntegral, SC_PERF_TEST_P(SCascadeTest, detectOnIntegral,
testing::Combine( testing::Combine(
testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), testing::Values(std::string("cv/cascadeandhog/cascades/inria_caltech-17.01.2013.xml"),
std::string("cv/cascadeandhog/cascades/sc_cvpr_2012_to_opencv_new_format.xml")),
testing::Values(std::string("cv/cascadeandhog/integrals.xml")))) testing::Values(std::string("cv/cascadeandhog/integrals.xml"))))
static std::string itoa(long i) static std::string itoa(long i)
@ -205,17 +209,12 @@ static std::string itoa(long i)
RUN_GPU(SCascadeTest, detectOnIntegral) RUN_GPU(SCascadeTest, detectOnIntegral)
{ {
cv::FileStorage fsi(perf::TestBase::getDataPath(GET_PARAM(1)), cv::FileStorage::READ); cv::Mat cpu = readImage ("cv/cascadeandhog/images/image_00000000_0.png");
ASSERT_TRUE(fsi.isOpened()); ASSERT_FALSE(cpu.empty());
cv::gpu::GpuMat hogluv(121 * 10, 161, CV_32SC1); cv::ICFPreprocessor preprocessor;
for (int i = 0; i < 10; ++i) cv::Mat test_res(cpu.rows / 4 * 10 + 1, cpu.cols / 4 + 1, CV_8UC1);
{ preprocessor.apply(cpu,test_res);
cv::Mat channel;
fsi[std::string("channel") + itoa(i)] >> channel;
cv::gpu::GpuMat gchannel(hogluv, cv::Rect(0, 121 * i, 161, 121));
gchannel.upload(channel);
}
cv::gpu::SCascade cascade; cv::gpu::SCascade cascade;
@ -227,6 +226,8 @@ RUN_GPU(SCascadeTest, detectOnIntegral)
cv::gpu::GpuMat objectBoxes(1, 10000 * sizeof(cv::gpu::SCascade::Detection), CV_8UC1), rois(cv::Size(640, 480), CV_8UC1); cv::gpu::GpuMat objectBoxes(1, 10000 * sizeof(cv::gpu::SCascade::Detection), CV_8UC1), rois(cv::Size(640, 480), CV_8UC1);
rois.setTo(1); rois.setTo(1);
cv::gpu::GpuMat hogluv(test_res);
cascade.detect(hogluv, rois, objectBoxes); cascade.detect(hogluv, rois, objectBoxes);
TEST_CYCLE() TEST_CYCLE()
@ -239,10 +240,11 @@ RUN_GPU(SCascadeTest, detectOnIntegral)
NO_CPU(SCascadeTest, detectOnIntegral) NO_CPU(SCascadeTest, detectOnIntegral)
PERF_TEST_P1(SCascadeTest, detectStream, SC_PERF_TEST_P(SCascadeTest, detectStream,
testing::Combine( testing::Combine(
testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), testing::Values(std::string("cv/cascadeandhog/cascades/inria_caltech-17.01.2013.xml"),
testing::Values(std::string("cv/cascadeandhog/bahnhof/image_00000000_0.png")))) std::string("cv/cascadeandhog/cascades/sc_cvpr_2012_to_opencv_new_format.xml")),
testing::Values(std::string("cv/cascadeandhog/images/image_00000000_0.png"))))
RUN_GPU(SCascadeTest, detectStream) RUN_GPU(SCascadeTest, detectStream)
{ {
@ -277,3 +279,5 @@ RUN_GPU(SCascadeTest, detectStream)
} }
NO_CPU(SCascadeTest, detectStream) NO_CPU(SCascadeTest, detectStream)
#undef SC_PERF_TEST_P

11
modules/gpu/src/cuda/icf-sc.cu
View File

@ -352,7 +352,7 @@ namespace icf {
{ {
#if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300 #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300
#pragma unroll #pragma unroll
// scan on shuffl functions // scan on shuffle functions
for (int i = 1; i < Policy::WARP; i *= 2) for (int i = 1; i < Policy::WARP; i *= 2)
{ {
const float n = __shfl_up(impact, i, Policy::WARP); const float n = __shfl_up(impact, i, Policy::WARP);
@ -459,7 +459,7 @@ __device void CascadeInvoker<Policy>::detect(Detection* objects, const uint ndet
const int y = blockIdx.y * blockDim.y + threadIdx.y; const int y = blockIdx.y * blockDim.y + threadIdx.y;
const int x = blockIdx.x; const int x = blockIdx.x;
// load Lavel // load Level
__shared__ Level level; __shared__ Level level;
// check POI // check POI
@ -501,11 +501,12 @@ __device void CascadeInvoker<Policy>::detect(Detection* objects, const uint ndet
float impact = leaves[(st + threadIdx.x) * 4 + lShift]; float impact = leaves[(st + threadIdx.x) * 4 + lShift];
PrefixSum<Policy>::apply(impact); PrefixSum<Policy>::apply(impact);
confidence += impact;
#if __CUDA_ARCH__ >= 120 #if __CUDA_ARCH__ >= 120
if(__any((confidence <= stages[(st + threadIdx.x)]))) st += 2048; if(__any((confidence + impact <= stages[(st + threadIdx.x)]))) st += 2048;
#endif #endif
impact = __shfl(impact, 31);
confidence += impact;
} }
if(!threadIdx.x && st == stEnd && ((confidence - FLT_EPSILON) >= 0)) if(!threadIdx.x && st == stEnd && ((confidence - FLT_EPSILON) >= 0))
@ -546,7 +547,7 @@ void CascadeInvoker<Policy>::operator()(const PtrStepSzb& roi, const PtrStepSzi&
soft_cascade<Policy, false><<<grid, Policy::block(), 0, stream>>>(inv, det, max_det, ctr, 0); soft_cascade<Policy, false><<<grid, Policy::block(), 0, stream>>>(inv, det, max_det, ctr, 0);
cudaSafeCall( cudaGetLastError()); cudaSafeCall( cudaGetLastError());
grid = dim3(fw, fh / Policy::STA_Y, scales - downscales); grid = dim3(fw, fh / Policy::STA_Y, 38 - downscales);
soft_cascade<Policy, true><<<grid, Policy::block(), 0, stream>>>(inv, det, max_det, ctr, downscales); soft_cascade<Policy, true><<<grid, Policy::block(), 0, stream>>>(inv, det, max_det, ctr, downscales);
if (!stream) if (!stream)

160
modules/gpu/src/softcascade.cpp
View File

@ -103,43 +103,44 @@ struct cv::gpu::SCascade::Fields
{ {
static const char *const SC_STAGE_TYPE = "stageType"; static const char *const SC_STAGE_TYPE = "stageType";
static const char *const SC_BOOST = "BOOST"; static const char *const SC_BOOST = "BOOST";
static const char *const SC_FEATURE_TYPE = "featureType"; static const char *const SC_FEATURE_TYPE = "featureType";
static const char *const SC_ICF = "ICF"; 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_FEATURE_FORMAT = "featureFormat";
static const char *const SC_SHRINKAGE = "shrinkage";
static const char *const SC_OCTAVES = "octaves";
static const char *const SC_OCT_SCALE = "scale";
static const char *const SC_OCT_WEAKS = "weaks";
static const char *const SC_TREES = "trees";
static const char *const SC_WEAK_THRESHOLD = "treeThreshold";
static const char *const SC_FEATURES = "features";
static const char *const SC_INTERNAL = "internalNodes";
static const char *const SC_LEAF = "leafValues";
static const char *const SC_F_CHANNEL = "channel";
static const char *const SC_F_RECT = "rect";
// only Ada Boost supported // only Ada Boost supported
std::string stageTypeStr = (string)root[SC_STAGE_TYPE]; std::string stageTypeStr = (string)root[SC_STAGE_TYPE];
CV_Assert(stageTypeStr == SC_BOOST); CV_Assert(stageTypeStr == SC_BOOST);
// only HOG-like integral channel features cupported // only HOG-like integral channel features supported
string featureTypeStr = (string)root[SC_FEATURE_TYPE]; string featureTypeStr = (string)root[SC_FEATURE_TYPE];
CV_Assert(featureTypeStr == SC_ICF); CV_Assert(featureTypeStr == SC_ICF);
static const char *const SC_ORIG_W = "width";
static const char *const SC_ORIG_H = "height";
int origWidth = (int)root[SC_ORIG_W]; int origWidth = (int)root[SC_ORIG_W];
int origHeight = (int)root[SC_ORIG_H]; int origHeight = (int)root[SC_ORIG_H];
static const char *const SC_OCTAVES = "octaves"; std::string fformat = (string)root[SC_FEATURE_FORMAT];
static const char *const SC_STAGES = "stages"; bool useBoxes = (fformat == "BOX");
static const char *const SC_FEATURES = "features";
static const char *const SC_WEEK = "weakClassifiers"; if(useBoxes)
static const char *const SC_INTERNAL = "internalNodes"; std::cout << "use boxes!!!";
static const char *const SC_LEAF = "leafValues";
static const char *const SC_OCT_SCALE = "scale"; ushort shrinkage = cv::saturate_cast<ushort>((int)root[SC_SHRINKAGE]);
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";
FileNode fn = root[SC_OCTAVES]; FileNode fn = root[SC_OCTAVES];
if (fn.empty()) return false; if (fn.empty()) return 0;
using namespace device::icf; using namespace device::icf;
@ -149,82 +150,105 @@ struct cv::gpu::SCascade::Fields
std::vector<float> vleaves; std::vector<float> vleaves;
FileNodeIterator it = fn.begin(), it_end = fn.end(); FileNodeIterator it = fn.begin(), it_end = fn.end();
int feature_offset = 0; for (ushort octIndex = 0; it != it_end; ++it, ++octIndex)
ushort octIndex = 0;
ushort shrinkage = 1;
for (; it != it_end; ++it)
{ {
FileNode fns = *it; FileNode fns = *it;
float scale = (float)fns[SC_OCT_SCALE]; float scale = powf(2.f,saturate_cast<float>((int)fns[SC_OCT_SCALE]));
std::cout << "octave scale " << scale << std::endl;
bool isUPOctave = scale >= 1; bool isUPOctave = scale >= 1;
ushort nstages = saturate_cast<ushort>((int)fns[SC_OCT_STAGES]); if (isUPOctave)
std::cout << "isUPOctave" << std::endl;
ushort nweaks = saturate_cast<ushort>((int)fns[SC_OCT_WEAKS]);
ushort2 size; ushort2 size;
size.x = cvRound(origWidth * scale); size.x = cvRound(origWidth * scale);
size.y = cvRound(origHeight * scale); size.y = cvRound(origHeight * scale);
shrinkage = saturate_cast<ushort>((int)fns[SC_OCT_SHRINKAGE]);
Octave octave(octIndex, nstages, shrinkage, size, scale); Octave octave(octIndex, nweaks, shrinkage, size, scale);
CV_Assert(octave.stages > 0); CV_Assert(octave.stages > 0);
voctaves.push_back(octave); voctaves.push_back(octave);
FileNode ffs = fns[SC_FEATURES]; FileNode ffs = fns[SC_FEATURES];
if (ffs.empty()) return false; if (ffs.empty()) return 0;
FileNodeIterator ftrs = ffs.begin(); std::vector<cv::Rect> feature_rects;
std::vector<int> feature_channels;
fns = fns[SC_STAGES]; FileNodeIterator ftrs = ffs.begin(), ftrs_end = ffs.end();
int feature_offset = 0;
for (; ftrs != ftrs_end; ++ftrs, ++feature_offset )
{
cv::FileNode ftn = (*ftrs)[SC_F_RECT];
cv::FileNodeIterator r_it = ftn.begin();
int x = (int)*(r_it++);
int y = (int)*(r_it++);
int w = (int)*(r_it++);
int h = (int)*(r_it++);
if (useBoxes)
{
if (isUPOctave)
{
w -= x;
h -= y;
}
}
else
{
if (!isUPOctave)
{
w += x;
h += y;
}
}
feature_rects.push_back(cv::Rect(x, y, w, h));
feature_channels.push_back((int)(*ftrs)[SC_F_CHANNEL]);
}
fns = fns[SC_TREES];
if (fn.empty()) return false; if (fn.empty()) return false;
// for each stage (~ decision tree with H = 2) // for each stage (~ decision tree with H = 2)
FileNodeIterator st = fns.begin(), st_end = fns.end(); FileNodeIterator st = fns.begin(), st_end = fns.end();
for (; st != st_end; ++st ) for (; st != st_end; ++st )
{ {
fns = *st; FileNode octfn = *st;
vstages.push_back((float)fns[SC_STAGE_THRESHOLD]); float threshold = (float)octfn[SC_WEAK_THRESHOLD];
vstages.push_back(threshold);
fns = fns[SC_WEEK]; FileNode intfns = octfn[SC_INTERNAL];
FileNodeIterator ftr = fns.begin(), ft_end = fns.end(); FileNodeIterator inIt = intfns.begin(), inIt_end = intfns.end();
for (; ftr != ft_end; ++ftr) for (; inIt != inIt_end;)
{ {
fns = (*ftr)[SC_INTERNAL]; inIt +=2;
FileNodeIterator inIt = fns.begin(), inIt_end = fns.end(); int featureIdx = (int)(*(inIt++));
for (; inIt != inIt_end;) // std::cout << " featureIdx " << featureIdx << " " << feature_rects[featureIdx] << std::endl;
{
// int feature = (int)(*(inIt +=2)) + feature_offset;
inIt +=3;
// extract feature, Todo:check it
unsigned int th = saturate_cast<unsigned int>((float)(*(inIt++)));
cv::FileNode ftn = (*ftrs)[SC_F_RECT];
cv::FileNodeIterator r_it = ftn.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++));
if (isUPOctave) float orig_threshold = (float)(*(inIt++));
{ unsigned int th = saturate_cast<unsigned int>((int)orig_threshold);
rect.z -= rect.x; // std::cout << "orig_threshold " << orig_threshold << " converted " << th << std::endl;
rect.w -= rect.y; cv::Rect& r = feature_rects[featureIdx];
} uchar4 rect;
rect.x = saturate_cast<uchar>(r.x);
rect.y = saturate_cast<uchar>(r.y);
rect.z = saturate_cast<uchar>(r.width);
rect.w = saturate_cast<uchar>(r.height);
unsigned int channel = saturate_cast<unsigned int>((int)(*ftrs)[SC_F_CHANNEL]); unsigned int channel = saturate_cast<unsigned int>(feature_channels[featureIdx]);
vnodes.push_back(Node(rect, channel, th)); vnodes.push_back(Node(rect, channel, th));
++ftrs; }
}
fns = (*ftr)[SC_LEAF]; intfns = octfn[SC_LEAF];
inIt = fns.begin(), inIt_end = fns.end(); inIt = intfns.begin(), inIt_end = intfns.end();
for (; inIt != inIt_end; ++inIt) for (; inIt != inIt_end; ++inIt)
vleaves.push_back((float)(*inIt)); {
vleaves.push_back((float)(*inIt));
} }
} }
std::cout << std::endl;
feature_offset += octave.stages * 3;
++octIndex;
} }
cv::Mat hoctaves(1, (int) (voctaves.size() * sizeof(Octave)), CV_8UC1, (uchar*)&(voctaves[0])); cv::Mat hoctaves(1, (int) (voctaves.size() * sizeof(Octave)), CV_8UC1, (uchar*)&(voctaves[0]));

119
modules/gpu/test/test_softcascade.cpp
View File

@ -51,20 +51,25 @@ using cv::gpu::GpuMat;
#if defined SHOW_DETECTIONS #if defined SHOW_DETECTIONS
# define SHOW(res) \ # define SHOW(res) \
cv::imshow(#res, result);\ cv::imshow(#res, res); \
cv::waitKey(0); cv::waitKey(0);
#else #else
# define SHOW(res) # define SHOW(res)
#endif #endif
static std::string path(std::string relative)
{
return cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/" + relative;
}
TEST(SCascadeTest, readCascade) TEST(SCascadeTest, readCascade)
{ {
std::string xml = cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/icf-template.xml"; std::string xml = path("cascades/inria_caltech-17.01.2013.xml");
cv::FileStorage fs(xml, cv::FileStorage::READ);
cv::gpu::SCascade cascade; cv::gpu::SCascade cascade;
cv::FileStorage fs(xml, cv::FileStorage::READ);
ASSERT_TRUE(fs.isOpened()); ASSERT_TRUE(fs.isOpened());
ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode()));
} }
@ -92,12 +97,6 @@ namespace
return rois[idx]; return rois[idx];
} }
std::string itoa(long i)
{
static char s[65];
sprintf(s, "%ld", i);
return std::string(s);
}
void print(std::ostream &out, const Detection& d) void print(std::ostream &out, const Detection& d)
{ {
@ -127,6 +126,13 @@ namespace
#endif #endif
} }
std::string itoa(long i)
{
static char s[65];
sprintf(s, "%ld", i);
return std::string(s);
}
#if defined SHOW_DETECTIONS #if defined SHOW_DETECTIONS
std::string getImageName(int level) std::string getImageName(int level)
{ {
@ -152,17 +158,20 @@ namespace
PARAM_TEST_CASE(SCascadeTestRoi, cv::gpu::DeviceInfo, std::string, std::string, int) PARAM_TEST_CASE(SCascadeTestRoi, cv::gpu::DeviceInfo, std::string, std::string, int)
{ {
virtual void SetUp()
{
cv::gpu::setDevice(GET_PARAM(0).deviceID());
}
}; };
GPU_TEST_P(SCascadeTestRoi, Detect) GPU_TEST_P(SCascadeTestRoi, Detect)
{ {
cv::gpu::setDevice(GET_PARAM(0).deviceID()); cv::Mat coloredCpu = cv::imread(path(GET_PARAM(2)));
cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() + GET_PARAM(2));
ASSERT_FALSE(coloredCpu.empty()); ASSERT_FALSE(coloredCpu.empty());
cv::gpu::SCascade cascade; cv::gpu::SCascade cascade;
cv::FileStorage fs(perf::TestBase::getDataPath(GET_PARAM(1)), cv::FileStorage::READ); cv::FileStorage fs(path(GET_PARAM(1)), cv::FileStorage::READ);
ASSERT_TRUE(fs.isOpened()); ASSERT_TRUE(fs.isOpened());
ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode()));
@ -204,21 +213,26 @@ GPU_TEST_P(SCascadeTestRoi, Detect)
INSTANTIATE_TEST_CASE_P(GPU_SoftCascade, SCascadeTestRoi, testing::Combine( INSTANTIATE_TEST_CASE_P(GPU_SoftCascade, SCascadeTestRoi, testing::Combine(
ALL_DEVICES, ALL_DEVICES,
testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), testing::Values(std::string("cascades/inria_caltech-17.01.2013.xml"),
testing::Values(std::string("../cv/cascadeandhog/bahnhof/image_00000000_0.png")), std::string("cascades/sc_cvpr_2012_to_opencv_new_format.xml")),
testing::Values(std::string("images/image_00000000_0.png")),
testing::Range(0, 5))); testing::Range(0, 5)));
struct SCascadeTestAll : testing::TestWithParam<cv::gpu::DeviceInfo> ////////////////////////////////////////
PARAM_TEST_CASE(SCascadeTestAll, cv::gpu::DeviceInfo, std::string)
{ {
std::string xml;
virtual void SetUp() virtual void SetUp()
{ {
cv::gpu::setDevice(GetParam().deviceID()); cv::gpu::setDevice(GET_PARAM(0).deviceID());
xml = path(GET_PARAM(1));
} }
}; };
GPU_TEST_P(SCascadeTestAll, detect) GPU_TEST_P(SCascadeTestAll, detect)
{ {
std::string xml = cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml";
cv::gpu::SCascade cascade; cv::gpu::SCascade cascade;
cv::FileStorage fs(xml, cv::FileStorage::READ); cv::FileStorage fs(xml, cv::FileStorage::READ);
@ -226,26 +240,36 @@ GPU_TEST_P(SCascadeTestAll, detect)
ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode()));
cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() cv::Mat coloredCpu = cv::imread(path("images/image_00000000_0.png"));
+ "../cv/cascadeandhog/bahnhof/image_00000000_0.png");
ASSERT_FALSE(coloredCpu.empty()); ASSERT_FALSE(coloredCpu.empty());
GpuMat colored(coloredCpu), objectBoxes, rois(colored.size(), CV_8UC1); GpuMat colored(coloredCpu), objectBoxes, rois(colored.size(), CV_8UC1);
rois.setTo(0); rois.setTo(1);
GpuMat sub(rois, cv::Rect(rois.cols / 4, rois.rows / 4,rois.cols / 2, rois.rows / 2));
sub.setTo(cv::Scalar::all(1));
cascade.detect(colored, rois, objectBoxes); cascade.detect(colored, rois, objectBoxes);
typedef cv::gpu::SCascade::Detection Detection; typedef cv::gpu::SCascade::Detection Detection;
cv::Mat detections(objectBoxes); cv::Mat dt(objectBoxes);
int a = *(detections.ptr<int>(0));
ASSERT_EQ(a, 2448);
Detection* dts = ((Detection*)dt.data) + 1;
int* count = dt.ptr<int>(0);
printTotal(std::cout, *count);
for (int i = 0; i < *count; ++i)
{
Detection d = dts[i];
print(std::cout, d);
cv::rectangle(coloredCpu, cv::Rect(d.x, d.y, d.w, d.h), cv::Scalar(255, 0, 0, 255), 1);
}
SHOW(coloredCpu);
// ASSERT_EQ(count, 2448);
} }
GPU_TEST_P(SCascadeTestAll, detectOnIntegral) GPU_TEST_P(SCascadeTestAll, detectOnIntegral)
{ {
std::string xml = cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml";
cv::gpu::SCascade cascade; cv::gpu::SCascade cascade;
cv::FileStorage fs(xml, cv::FileStorage::READ); cv::FileStorage fs(xml, cv::FileStorage::READ);
@ -253,18 +277,11 @@ GPU_TEST_P(SCascadeTestAll, detectOnIntegral)
ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode()));
std::string intPath = cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/integrals.xml"; cv::Mat coloredCpu = cv::imread(path("images/image_00000000_0.png"));
cv::FileStorage fsi(intPath, cv::FileStorage::READ); cv::ICFPreprocessor preprocessor;
ASSERT_TRUE(fsi.isOpened()); cv::Mat integrals(coloredCpu.rows / 4 * 10 + 1, coloredCpu.cols / 4 + 1, CV_8UC1);
preprocessor.apply(coloredCpu, integrals);
GpuMat hogluv(121 * 10, 161, CV_32SC1); GpuMat hogluv(integrals);
for (int i = 0; i < 10; ++i)
{
cv::Mat channel;
fsi[std::string("channel") + itoa(i)] >> channel;
GpuMat gchannel(hogluv, cv::Rect(0, 121 * i, 161, 121));
gchannel.upload(channel);
}
GpuMat objectBoxes(1, 100000, CV_8UC1), rois(cv::Size(640, 480), CV_8UC1); GpuMat objectBoxes(1, 100000, CV_8UC1), rois(cv::Size(640, 480), CV_8UC1);
rois.setTo(1); rois.setTo(1);
@ -272,15 +289,14 @@ GPU_TEST_P(SCascadeTestAll, detectOnIntegral)
objectBoxes.setTo(0); objectBoxes.setTo(0);
cascade.detect(hogluv, rois, objectBoxes); cascade.detect(hogluv, rois, objectBoxes);
typedef cv::gpu::SCascade::Detection Detection; // typedef cv::gpu::SCascade::Detection Detection;
cv::Mat detections(objectBoxes); // cv::Mat detections(objectBoxes);
int a = *(detections.ptr<int>(0)); // int a = *(detections.ptr<int>(0));
ASSERT_EQ(a, 1024); // ASSERT_EQ(a, 1024);
} }
GPU_TEST_P(SCascadeTestAll, detectStream) GPU_TEST_P(SCascadeTestAll, detectStream)
{ {
std::string xml = cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml";
cv::gpu::SCascade cascade; cv::gpu::SCascade cascade;
cv::FileStorage fs(xml, cv::FileStorage::READ); cv::FileStorage fs(xml, cv::FileStorage::READ);
@ -288,8 +304,7 @@ GPU_TEST_P(SCascadeTestAll, detectStream)
ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode()));
cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() cv::Mat coloredCpu = cv::imread(path("images/image_00000000_0.png"));
+ "../cv/cascadeandhog/bahnhof/image_00000000_0.png");
ASSERT_FALSE(coloredCpu.empty()); ASSERT_FALSE(coloredCpu.empty());
GpuMat colored(coloredCpu), objectBoxes(1, 100000, CV_8UC1), rois(colored.size(), CV_8UC1); GpuMat colored(coloredCpu), objectBoxes(1, 100000, CV_8UC1), rois(colored.size(), CV_8UC1);
@ -303,12 +318,14 @@ GPU_TEST_P(SCascadeTestAll, detectStream)
cascade.detect(colored, rois, objectBoxes, s); cascade.detect(colored, rois, objectBoxes, s);
s.waitForCompletion(); s.waitForCompletion();
typedef cv::gpu::SCascade::Detection Detection; // typedef cv::gpu::SCascade::Detection Detection;
cv::Mat detections(objectBoxes); // cv::Mat detections(objectBoxes);
int a = *(detections.ptr<int>(0)); // int a = *(detections.ptr<int>(0));
ASSERT_EQ(a, 2448); // ASSERT_EQ(a, 2448);
} }
INSTANTIATE_TEST_CASE_P(GPU_SoftCascade, SCascadeTestAll, ALL_DEVICES); INSTANTIATE_TEST_CASE_P(GPU_SoftCascade, SCascadeTestAll, testing::Combine( ALL_DEVICES,
testing::Values(std::string("cascades/inria_caltech-17.01.2013.xml"),
std::string("cascades/sc_cvpr_2012_to_opencv_new_format.xml"))));
#endif #endif

9
modules/objdetect/include/opencv2/objdetect/objdetect.hpp
View File

@ -489,6 +489,15 @@ protected:
}; };
class CV_EXPORTS_W ICFPreprocessor
{
public:
CV_WRAP ICFPreprocessor();
CV_WRAP void apply(cv::InputArray _frame, cv::OutputArray _integrals) const;
protected:
enum {BINS = 10};
};
// Implementation of soft (stageless) cascaded detector. // Implementation of soft (stageless) cascaded detector.
class CV_EXPORTS_W SCascade : public Algorithm class CV_EXPORTS_W SCascade : public Algorithm
{ {

3
modules/objdetect/src/softcascade.cpp
View File

@ -41,7 +41,6 @@
//M*/ //M*/
#include "precomp.hpp" #include "precomp.hpp"
#include <iostream>
namespace { namespace {
@ -365,7 +364,7 @@ struct cv::SCascade::Fields
std::string fformat = (string)root[FEATURE_FORMAT]; std::string fformat = (string)root[FEATURE_FORMAT];
bool useBoxes = (fformat == "BOX"); bool useBoxes = (fformat == "BOX");
// only HOG-like integral channel features cupported // only HOG-like integral channel features supported
string featureTypeStr = (string)root[SC_FEATURE_TYPE]; string featureTypeStr = (string)root[SC_FEATURE_TYPE];
CV_Assert(featureTypeStr == SC_ICF); CV_Assert(featureTypeStr == SC_ICF);