From 8d9c9c269080860481e84426ebe14669e6920bd5 Mon Sep 17 00:00:00 2001 From: "marina.kolpakova" Date: Wed, 23 Jan 2013 14:57:51 +0400 Subject: [PATCH 1/2] integrate new cascade format to GPU soft cascade implementation --- modules/gpu/include/opencv2/gpu/gpu.hpp | 8 +- modules/gpu/perf/perf_softcascade.cpp | 56 +++--- modules/gpu/src/cuda/icf-sc.cu | 11 +- modules/gpu/src/softcascade.cpp | 160 ++++++++++-------- modules/gpu/test/test_softcascade.cpp | 119 +++++++------ .../include/opencv2/objdetect/objdetect.hpp | 9 + modules/objdetect/src/softcascade.cpp | 3 +- 7 files changed, 210 insertions(+), 156 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index c6ce2faff..ed3232622 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1556,7 +1556,7 @@ protected: ChannelsProcessor(); }; -// Implementation of soft (stageless) cascaded detector. +// Implementation of soft (stage-less) cascaded detector. class CV_EXPORTS SCascade : public Algorithm { public: @@ -1577,8 +1577,8 @@ public: 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 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 applied. // Param scales is a number of scales from minScale to maxScale. // Param flags is an extra tuning flags. SCascade(const double minScale = 0.4, const double maxScale = 5., const int scales = 55, @@ -1595,7 +1595,7 @@ public: // Load cascade config. 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 rois is a regions of interests mask generated by genRoi. // Only the objects that fall into one of the regions will be returned. diff --git a/modules/gpu/perf/perf_softcascade.cpp b/modules/gpu/perf/perf_softcascade.cpp index 32e41a432..97ee9de87 100644 --- a/modules/gpu/perf/perf_softcascade.cpp +++ b/modules/gpu/perf/perf_softcascade.cpp @@ -1,6 +1,6 @@ #include "perf_precomp.hpp" -#define PERF_TEST_P1(fixture, name, params) \ +#define SC_PERF_TEST_P(fixture, name, params) \ class fixture##_##name : public fixture {\ public:\ fixture##_##name() {}\ @@ -28,7 +28,7 @@ namespace { bool operator()(const cv::gpu::SCascade::Detection& a, 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.w != b.w) return a.w < b.w; else return a.h < b.h; @@ -52,10 +52,11 @@ namespace { typedef std::tr1::tuple fixture_t; typedef perf::TestBaseWithParam SCascadeTest; -PERF_TEST_P1(SCascadeTest, detect, +SC_PERF_TEST_P(SCascadeTest, detect, testing::Combine( - testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), - testing::Values(std::string("cv/cascadeandhog/bahnhof/image_00000000_0.png")))) + 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/images/image_00000000_0.png")))) RUN_GPU(SCascadeTest, detect) { @@ -108,10 +109,11 @@ static cv::Rect getFromTable(int idx) typedef std::tr1::tuple roi_fixture_t; typedef perf::TestBaseWithParam SCascadeTestRoi; -PERF_TEST_P1(SCascadeTestRoi, detectInRoi, +SC_PERF_TEST_P(SCascadeTestRoi, detectInRoi, testing::Combine( - testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), - testing::Values(std::string("cv/cascadeandhog/bahnhof/image_00000000_0.png")), + 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/images/image_00000000_0.png")), testing::Range(0, 5))) RUN_GPU(SCascadeTestRoi, detectInRoi) @@ -152,10 +154,11 @@ RUN_GPU(SCascadeTestRoi, detectInRoi) NO_CPU(SCascadeTestRoi, detectInRoi) -PERF_TEST_P1(SCascadeTestRoi, detectEachRoi, +SC_PERF_TEST_P(SCascadeTestRoi, detectEachRoi, testing::Combine( - testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), - testing::Values(std::string("cv/cascadeandhog/bahnhof/image_00000000_0.png")), + 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/images/image_00000000_0.png")), testing::Range(0, 10))) RUN_GPU(SCascadeTestRoi, detectEachRoi) @@ -191,9 +194,10 @@ RUN_GPU(SCascadeTestRoi, detectEachRoi) NO_CPU(SCascadeTestRoi, detectEachRoi) -PERF_TEST_P1(SCascadeTest, detectOnIntegral, +SC_PERF_TEST_P(SCascadeTest, detectOnIntegral, 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")))) static std::string itoa(long i) @@ -205,17 +209,12 @@ static std::string itoa(long i) RUN_GPU(SCascadeTest, detectOnIntegral) { - cv::FileStorage fsi(perf::TestBase::getDataPath(GET_PARAM(1)), cv::FileStorage::READ); - ASSERT_TRUE(fsi.isOpened()); + cv::Mat cpu = readImage ("cv/cascadeandhog/images/image_00000000_0.png"); + ASSERT_FALSE(cpu.empty()); - cv::gpu::GpuMat hogluv(121 * 10, 161, CV_32SC1); - for (int i = 0; i < 10; ++i) - { - 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::ICFPreprocessor preprocessor; + cv::Mat test_res(cpu.rows / 4 * 10 + 1, cpu.cols / 4 + 1, CV_8UC1); + preprocessor.apply(cpu,test_res); 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); rois.setTo(1); + cv::gpu::GpuMat hogluv(test_res); + cascade.detect(hogluv, rois, objectBoxes); TEST_CYCLE() @@ -239,10 +240,11 @@ RUN_GPU(SCascadeTest, detectOnIntegral) NO_CPU(SCascadeTest, detectOnIntegral) -PERF_TEST_P1(SCascadeTest, detectStream, +SC_PERF_TEST_P(SCascadeTest, detectStream, testing::Combine( - testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), - testing::Values(std::string("cv/cascadeandhog/bahnhof/image_00000000_0.png")))) + 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/images/image_00000000_0.png")))) RUN_GPU(SCascadeTest, detectStream) { @@ -277,3 +279,5 @@ RUN_GPU(SCascadeTest, detectStream) } NO_CPU(SCascadeTest, detectStream) + +#undef SC_PERF_TEST_P \ No newline at end of file diff --git a/modules/gpu/src/cuda/icf-sc.cu b/modules/gpu/src/cuda/icf-sc.cu index aa695372e..1894cd45e 100644 --- a/modules/gpu/src/cuda/icf-sc.cu +++ b/modules/gpu/src/cuda/icf-sc.cu @@ -352,7 +352,7 @@ namespace icf { { #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 300 #pragma unroll - // scan on shuffl functions + // scan on shuffle functions for (int i = 1; i < Policy::WARP; i *= 2) { const float n = __shfl_up(impact, i, Policy::WARP); @@ -459,7 +459,7 @@ __device void CascadeInvoker::detect(Detection* objects, const uint ndet const int y = blockIdx.y * blockDim.y + threadIdx.y; const int x = blockIdx.x; - // load Lavel + // load Level __shared__ Level level; // check POI @@ -501,11 +501,12 @@ __device void CascadeInvoker::detect(Detection* objects, const uint ndet float impact = leaves[(st + threadIdx.x) * 4 + lShift]; PrefixSum::apply(impact); - confidence += impact; #if __CUDA_ARCH__ >= 120 - if(__any((confidence <= stages[(st + threadIdx.x)]))) st += 2048; + if(__any((confidence + impact <= stages[(st + threadIdx.x)]))) st += 2048; #endif + impact = __shfl(impact, 31); + confidence += impact; } if(!threadIdx.x && st == stEnd && ((confidence - FLT_EPSILON) >= 0)) @@ -546,7 +547,7 @@ void CascadeInvoker::operator()(const PtrStepSzb& roi, const PtrStepSzi& soft_cascade<<>>(inv, det, max_det, ctr, 0); cudaSafeCall( cudaGetLastError()); - grid = dim3(fw, fh / Policy::STA_Y, scales - downscales); + grid = dim3(fw, fh / Policy::STA_Y, 38 - downscales); soft_cascade<<>>(inv, det, max_det, ctr, downscales); if (!stream) diff --git a/modules/gpu/src/softcascade.cpp b/modules/gpu/src/softcascade.cpp index 561d4e44e..27c44d0ff 100644 --- a/modules/gpu/src/softcascade.cpp +++ b/modules/gpu/src/softcascade.cpp @@ -103,43 +103,44 @@ struct cv::gpu::SCascade::Fields { 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_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 std::string stageTypeStr = (string)root[SC_STAGE_TYPE]; 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]; 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 origHeight = (int)root[SC_ORIG_H]; - static const char *const SC_OCTAVES = "octaves"; - static const char *const SC_STAGES = "stages"; - static const char *const SC_FEATURES = "features"; + std::string fformat = (string)root[SC_FEATURE_FORMAT]; + bool useBoxes = (fformat == "BOX"); - static const char *const SC_WEEK = "weakClassifiers"; - static const char *const SC_INTERNAL = "internalNodes"; - static const char *const SC_LEAF = "leafValues"; + if(useBoxes) + std::cout << "use boxes!!!"; - 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"; + ushort shrinkage = cv::saturate_cast((int)root[SC_SHRINKAGE]); FileNode fn = root[SC_OCTAVES]; - if (fn.empty()) return false; + if (fn.empty()) return 0; using namespace device::icf; @@ -149,82 +150,105 @@ struct cv::gpu::SCascade::Fields std::vector vleaves; FileNodeIterator it = fn.begin(), it_end = fn.end(); - int feature_offset = 0; - ushort octIndex = 0; - ushort shrinkage = 1; - - for (; it != it_end; ++it) + for (ushort octIndex = 0; it != it_end; ++it, ++octIndex) { FileNode fns = *it; - float scale = (float)fns[SC_OCT_SCALE]; + float scale = powf(2.f,saturate_cast((int)fns[SC_OCT_SCALE])); + std::cout << "octave scale " << scale << std::endl; bool isUPOctave = scale >= 1; - ushort nstages = saturate_cast((int)fns[SC_OCT_STAGES]); + if (isUPOctave) + std::cout << "isUPOctave" << std::endl; + + ushort nweaks = saturate_cast((int)fns[SC_OCT_WEAKS]); + ushort2 size; size.x = cvRound(origWidth * scale); size.y = cvRound(origHeight * scale); - shrinkage = saturate_cast((int)fns[SC_OCT_SHRINKAGE]); - Octave octave(octIndex, nstages, shrinkage, size, scale); + Octave octave(octIndex, nweaks, shrinkage, size, scale); CV_Assert(octave.stages > 0); voctaves.push_back(octave); FileNode ffs = fns[SC_FEATURES]; - if (ffs.empty()) return false; + if (ffs.empty()) return 0; - FileNodeIterator ftrs = ffs.begin(); + std::vector feature_rects; + std::vector 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; // 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)fns[SC_STAGE_THRESHOLD]); + FileNode octfn = *st; + float threshold = (float)octfn[SC_WEAK_THRESHOLD]; + vstages.push_back(threshold); - fns = fns[SC_WEEK]; - FileNodeIterator ftr = fns.begin(), ft_end = fns.end(); - for (; ftr != ft_end; ++ftr) + FileNode intfns = octfn[SC_INTERNAL]; + FileNodeIterator inIt = intfns.begin(), inIt_end = intfns.end(); + for (; inIt != inIt_end;) { - fns = (*ftr)[SC_INTERNAL]; - FileNodeIterator inIt = fns.begin(), inIt_end = fns.end(); - for (; inIt != inIt_end;) - { - // int feature = (int)(*(inIt +=2)) + feature_offset; - inIt +=3; - // extract feature, Todo:check it - unsigned int th = saturate_cast((float)(*(inIt++))); - cv::FileNode ftn = (*ftrs)[SC_F_RECT]; - cv::FileNodeIterator r_it = ftn.begin(); - uchar4 rect; - rect.x = saturate_cast((int)*(r_it++)); - rect.y = saturate_cast((int)*(r_it++)); - rect.z = saturate_cast((int)*(r_it++)); - rect.w = saturate_cast((int)*(r_it++)); + inIt +=2; + int featureIdx = (int)(*(inIt++)); + // std::cout << " featureIdx " << featureIdx << " " << feature_rects[featureIdx] << std::endl; - if (isUPOctave) - { - rect.z -= rect.x; - rect.w -= rect.y; - } + float orig_threshold = (float)(*(inIt++)); + unsigned int th = saturate_cast((int)orig_threshold); + // std::cout << "orig_threshold " << orig_threshold << " converted " << th << std::endl; + cv::Rect& r = feature_rects[featureIdx]; + uchar4 rect; + rect.x = saturate_cast(r.x); + rect.y = saturate_cast(r.y); + rect.z = saturate_cast(r.width); + rect.w = saturate_cast(r.height); - unsigned int channel = saturate_cast((int)(*ftrs)[SC_F_CHANNEL]); - vnodes.push_back(Node(rect, channel, th)); - ++ftrs; - } + unsigned int channel = saturate_cast(feature_channels[featureIdx]); + vnodes.push_back(Node(rect, channel, th)); + } - fns = (*ftr)[SC_LEAF]; - inIt = fns.begin(), inIt_end = fns.end(); - for (; inIt != inIt_end; ++inIt) - vleaves.push_back((float)(*inIt)); + intfns = octfn[SC_LEAF]; + inIt = intfns.begin(), inIt_end = intfns.end(); + for (; inIt != inIt_end; ++inIt) + { + vleaves.push_back((float)(*inIt)); } } - - feature_offset += octave.stages * 3; - ++octIndex; + std::cout << std::endl; } cv::Mat hoctaves(1, (int) (voctaves.size() * sizeof(Octave)), CV_8UC1, (uchar*)&(voctaves[0])); diff --git a/modules/gpu/test/test_softcascade.cpp b/modules/gpu/test/test_softcascade.cpp index 9cc1a5e39..5b897a971 100644 --- a/modules/gpu/test/test_softcascade.cpp +++ b/modules/gpu/test/test_softcascade.cpp @@ -51,20 +51,25 @@ using cv::gpu::GpuMat; #if defined SHOW_DETECTIONS # define SHOW(res) \ - cv::imshow(#res, result);\ + cv::imshow(#res, res); \ cv::waitKey(0); #else # define SHOW(res) #endif +static std::string path(std::string relative) +{ + return cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/" + relative; +} + 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::FileStorage fs(xml, cv::FileStorage::READ); ASSERT_TRUE(fs.isOpened()); - ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); } @@ -92,12 +97,6 @@ namespace 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) { @@ -127,6 +126,13 @@ namespace #endif } + std::string itoa(long i) + { + static char s[65]; + sprintf(s, "%ld", i); + return std::string(s); + } + #if defined SHOW_DETECTIONS std::string getImageName(int level) { @@ -152,17 +158,20 @@ namespace 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) { - cv::gpu::setDevice(GET_PARAM(0).deviceID()); - cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() + GET_PARAM(2)); + cv::Mat coloredCpu = cv::imread(path(GET_PARAM(2))); ASSERT_FALSE(coloredCpu.empty()); 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(cascade.load(fs.getFirstTopLevelNode())); @@ -204,21 +213,26 @@ GPU_TEST_P(SCascadeTestRoi, Detect) INSTANTIATE_TEST_CASE_P(GPU_SoftCascade, SCascadeTestRoi, testing::Combine( ALL_DEVICES, - testing::Values(std::string("cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml")), - testing::Values(std::string("../cv/cascadeandhog/bahnhof/image_00000000_0.png")), + testing::Values(std::string("cascades/inria_caltech-17.01.2013.xml"), + std::string("cascades/sc_cvpr_2012_to_opencv_new_format.xml")), + testing::Values(std::string("images/image_00000000_0.png")), testing::Range(0, 5))); -struct SCascadeTestAll : testing::TestWithParam +//////////////////////////////////////// +PARAM_TEST_CASE(SCascadeTestAll, cv::gpu::DeviceInfo, std::string) { + + std::string xml; + 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) { - std::string xml = cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml"; cv::gpu::SCascade cascade; cv::FileStorage fs(xml, cv::FileStorage::READ); @@ -226,26 +240,36 @@ GPU_TEST_P(SCascadeTestAll, detect) ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); - cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() - + "../cv/cascadeandhog/bahnhof/image_00000000_0.png"); + cv::Mat coloredCpu = cv::imread(path("images/image_00000000_0.png")); ASSERT_FALSE(coloredCpu.empty()); GpuMat colored(coloredCpu), objectBoxes, rois(colored.size(), CV_8UC1); - rois.setTo(0); - GpuMat sub(rois, cv::Rect(rois.cols / 4, rois.rows / 4,rois.cols / 2, rois.rows / 2)); - sub.setTo(cv::Scalar::all(1)); + rois.setTo(1); cascade.detect(colored, rois, objectBoxes); typedef cv::gpu::SCascade::Detection Detection; - cv::Mat detections(objectBoxes); - int a = *(detections.ptr(0)); - ASSERT_EQ(a, 2448); + cv::Mat dt(objectBoxes); + + + Detection* dts = ((Detection*)dt.data) + 1; + int* count = dt.ptr(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) { - std::string xml = cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/sc_cvpr_2012_to_opencv.xml"; cv::gpu::SCascade cascade; cv::FileStorage fs(xml, cv::FileStorage::READ); @@ -253,18 +277,11 @@ GPU_TEST_P(SCascadeTestAll, detectOnIntegral) ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); - std::string intPath = cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/integrals.xml"; - cv::FileStorage fsi(intPath, cv::FileStorage::READ); - ASSERT_TRUE(fsi.isOpened()); - - GpuMat hogluv(121 * 10, 161, CV_32SC1); - 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); - } + cv::Mat coloredCpu = cv::imread(path("images/image_00000000_0.png")); + cv::ICFPreprocessor preprocessor; + cv::Mat integrals(coloredCpu.rows / 4 * 10 + 1, coloredCpu.cols / 4 + 1, CV_8UC1); + preprocessor.apply(coloredCpu, integrals); + GpuMat hogluv(integrals); GpuMat objectBoxes(1, 100000, CV_8UC1), rois(cv::Size(640, 480), CV_8UC1); rois.setTo(1); @@ -272,15 +289,14 @@ GPU_TEST_P(SCascadeTestAll, detectOnIntegral) objectBoxes.setTo(0); cascade.detect(hogluv, rois, objectBoxes); - typedef cv::gpu::SCascade::Detection Detection; - cv::Mat detections(objectBoxes); - int a = *(detections.ptr(0)); - ASSERT_EQ(a, 1024); +// typedef cv::gpu::SCascade::Detection Detection; +// cv::Mat detections(objectBoxes); +// int a = *(detections.ptr(0)); +// ASSERT_EQ(a, 1024); } 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::FileStorage fs(xml, cv::FileStorage::READ); @@ -288,8 +304,7 @@ GPU_TEST_P(SCascadeTestAll, detectStream) ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); - cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() - + "../cv/cascadeandhog/bahnhof/image_00000000_0.png"); + cv::Mat coloredCpu = cv::imread(path("images/image_00000000_0.png")); ASSERT_FALSE(coloredCpu.empty()); 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); s.waitForCompletion(); - typedef cv::gpu::SCascade::Detection Detection; - cv::Mat detections(objectBoxes); - int a = *(detections.ptr(0)); - ASSERT_EQ(a, 2448); + // typedef cv::gpu::SCascade::Detection Detection; + // cv::Mat detections(objectBoxes); + // int a = *(detections.ptr(0)); + // 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 diff --git a/modules/objdetect/include/opencv2/objdetect/objdetect.hpp b/modules/objdetect/include/opencv2/objdetect/objdetect.hpp index 986c4f6d8..95300c48a 100644 --- a/modules/objdetect/include/opencv2/objdetect/objdetect.hpp +++ b/modules/objdetect/include/opencv2/objdetect/objdetect.hpp @@ -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. class CV_EXPORTS_W SCascade : public Algorithm { diff --git a/modules/objdetect/src/softcascade.cpp b/modules/objdetect/src/softcascade.cpp index 8ce247c93..9893f417b 100644 --- a/modules/objdetect/src/softcascade.cpp +++ b/modules/objdetect/src/softcascade.cpp @@ -41,7 +41,6 @@ //M*/ #include "precomp.hpp" -#include namespace { @@ -365,7 +364,7 @@ struct cv::SCascade::Fields std::string fformat = (string)root[FEATURE_FORMAT]; 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]; CV_Assert(featureTypeStr == SC_ICF); From cb329400df0a3ed8b96397a74a967dc5add2615d Mon Sep 17 00:00:00 2001 From: "marina.kolpakova" Date: Wed, 30 Jan 2013 08:18:05 +0400 Subject: [PATCH 2/2] fix tests --- modules/gpu/perf/perf_softcascade.cpp | 51 +-------------- modules/gpu/src/cuda/icf-sc.cu | 2 +- modules/gpu/src/softcascade.cpp | 16 +---- modules/gpu/test/test_softcascade.cpp | 64 ++++++++----------- .../include/opencv2/objdetect/objdetect.hpp | 10 --- 5 files changed, 29 insertions(+), 114 deletions(-) diff --git a/modules/gpu/perf/perf_softcascade.cpp b/modules/gpu/perf/perf_softcascade.cpp index 97ee9de87..6cb3c6356 100644 --- a/modules/gpu/perf/perf_softcascade.cpp +++ b/modules/gpu/perf/perf_softcascade.cpp @@ -194,52 +194,6 @@ RUN_GPU(SCascadeTestRoi, detectEachRoi) NO_CPU(SCascadeTestRoi, detectEachRoi) -SC_PERF_TEST_P(SCascadeTest, detectOnIntegral, - testing::Combine( - 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")))) - -static std::string itoa(long i) -{ - static char s[65]; - sprintf(s, "%ld", i); - return std::string(s); -} - -RUN_GPU(SCascadeTest, detectOnIntegral) -{ - cv::Mat cpu = readImage ("cv/cascadeandhog/images/image_00000000_0.png"); - ASSERT_FALSE(cpu.empty()); - - cv::ICFPreprocessor preprocessor; - cv::Mat test_res(cpu.rows / 4 * 10 + 1, cpu.cols / 4 + 1, CV_8UC1); - preprocessor.apply(cpu,test_res); - - cv::gpu::SCascade cascade; - - cv::FileStorage fs(perf::TestBase::getDataPath(GET_PARAM(0)), cv::FileStorage::READ); - ASSERT_TRUE(fs.isOpened()); - - 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); - rois.setTo(1); - - cv::gpu::GpuMat hogluv(test_res); - - cascade.detect(hogluv, rois, objectBoxes); - - TEST_CYCLE() - { - cascade.detect(hogluv, rois, objectBoxes); - } - - SANITY_CHECK(sortDetections(objectBoxes)); -} - -NO_CPU(SCascadeTest, detectOnIntegral) - SC_PERF_TEST_P(SCascadeTest, detectStream, testing::Combine( testing::Values(std::string("cv/cascadeandhog/cascades/inria_caltech-17.01.2013.xml"), @@ -271,10 +225,7 @@ RUN_GPU(SCascadeTest, detectStream) cascade.detect(colored, rois, objectBoxes, s); } -#ifdef HAVE_CUDA - cudaDeviceSynchronize(); -#endif - + s.waitForCompletion(); SANITY_CHECK(sortDetections(objectBoxes)); } diff --git a/modules/gpu/src/cuda/icf-sc.cu b/modules/gpu/src/cuda/icf-sc.cu index 1894cd45e..d64c7a807 100644 --- a/modules/gpu/src/cuda/icf-sc.cu +++ b/modules/gpu/src/cuda/icf-sc.cu @@ -547,7 +547,7 @@ void CascadeInvoker::operator()(const PtrStepSzb& roi, const PtrStepSzi& soft_cascade<<>>(inv, det, max_det, ctr, 0); cudaSafeCall( cudaGetLastError()); - grid = dim3(fw, fh / Policy::STA_Y, 38 - downscales); + grid = dim3(fw, fh / Policy::STA_Y, min(38, scales) - downscales); soft_cascade<<>>(inv, det, max_det, ctr, downscales); if (!stream) diff --git a/modules/gpu/src/softcascade.cpp b/modules/gpu/src/softcascade.cpp index 27c44d0ff..88402092e 100644 --- a/modules/gpu/src/softcascade.cpp +++ b/modules/gpu/src/softcascade.cpp @@ -133,10 +133,6 @@ struct cv::gpu::SCascade::Fields std::string fformat = (string)root[SC_FEATURE_FORMAT]; bool useBoxes = (fformat == "BOX"); - - if(useBoxes) - std::cout << "use boxes!!!"; - ushort shrinkage = cv::saturate_cast((int)root[SC_SHRINKAGE]); FileNode fn = root[SC_OCTAVES]; @@ -154,13 +150,8 @@ struct cv::gpu::SCascade::Fields { FileNode fns = *it; float scale = powf(2.f,saturate_cast((int)fns[SC_OCT_SCALE])); - std::cout << "octave scale " << scale << std::endl; - bool isUPOctave = scale >= 1; - if (isUPOctave) - std::cout << "isUPOctave" << std::endl; - ushort nweaks = saturate_cast((int)fns[SC_OCT_WEAKS]); ushort2 size; @@ -225,11 +216,9 @@ struct cv::gpu::SCascade::Fields { inIt +=2; int featureIdx = (int)(*(inIt++)); - // std::cout << " featureIdx " << featureIdx << " " << feature_rects[featureIdx] << std::endl; float orig_threshold = (float)(*(inIt++)); unsigned int th = saturate_cast((int)orig_threshold); - // std::cout << "orig_threshold " << orig_threshold << " converted " << th << std::endl; cv::Rect& r = feature_rects[featureIdx]; uchar4 rect; rect.x = saturate_cast(r.x); @@ -248,7 +237,6 @@ struct cv::gpu::SCascade::Fields vleaves.push_back((float)(*inIt)); } } - std::cout << std::endl; } cv::Mat hoctaves(1, (int) (voctaves.size() * sizeof(Octave)), CV_8UC1, (uchar*)&(voctaves[0])); @@ -424,7 +412,7 @@ public: // 160x120x10 GpuMat shrunk; - // temporial mat for integrall + // temporal mat for integral GpuMat integralBuffer; // 161x121x10 @@ -591,7 +579,7 @@ private: cv::gpu::cartToPolar(dfdx, dfdy, mag, ang, true, s); - // normolize magnitude to uchar interval and angles to 6 bins + // normalize 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)); diff --git a/modules/gpu/test/test_softcascade.cpp b/modules/gpu/test/test_softcascade.cpp index 5b897a971..c08dc06c8 100644 --- a/modules/gpu/test/test_softcascade.cpp +++ b/modules/gpu/test/test_softcascade.cpp @@ -219,15 +219,30 @@ INSTANTIATE_TEST_CASE_P(GPU_SoftCascade, SCascadeTestRoi, testing::Combine( testing::Range(0, 5))); //////////////////////////////////////// -PARAM_TEST_CASE(SCascadeTestAll, cv::gpu::DeviceInfo, std::string) + +namespace { + +struct Fixture +{ + std::string path; + int expected; + + Fixture(){} + Fixture(std::string p, int e): path(p), expected(e) {} +}; +} + +PARAM_TEST_CASE(SCascadeTestAll, cv::gpu::DeviceInfo, Fixture) { std::string xml; + int expected; virtual void SetUp() { cv::gpu::setDevice(GET_PARAM(0).deviceID()); - xml = path(GET_PARAM(1)); + xml = path(GET_PARAM(1).path); + expected = GET_PARAM(1).expected; } }; @@ -265,34 +280,7 @@ GPU_TEST_P(SCascadeTestAll, detect) } SHOW(coloredCpu); - // ASSERT_EQ(count, 2448); -} - -GPU_TEST_P(SCascadeTestAll, detectOnIntegral) -{ - cv::gpu::SCascade cascade; - - cv::FileStorage fs(xml, cv::FileStorage::READ); - ASSERT_TRUE(fs.isOpened()); - - ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); - - cv::Mat coloredCpu = cv::imread(path("images/image_00000000_0.png")); - cv::ICFPreprocessor preprocessor; - cv::Mat integrals(coloredCpu.rows / 4 * 10 + 1, coloredCpu.cols / 4 + 1, CV_8UC1); - preprocessor.apply(coloredCpu, integrals); - GpuMat hogluv(integrals); - - GpuMat objectBoxes(1, 100000, CV_8UC1), rois(cv::Size(640, 480), CV_8UC1); - rois.setTo(1); - - objectBoxes.setTo(0); - cascade.detect(hogluv, rois, objectBoxes); - -// typedef cv::gpu::SCascade::Detection Detection; -// cv::Mat detections(objectBoxes); -// int a = *(detections.ptr(0)); -// ASSERT_EQ(a, 1024); + ASSERT_EQ(*count, expected); } GPU_TEST_P(SCascadeTestAll, detectStream) @@ -308,9 +296,7 @@ GPU_TEST_P(SCascadeTestAll, detectStream) ASSERT_FALSE(coloredCpu.empty()); GpuMat colored(coloredCpu), objectBoxes(1, 100000, CV_8UC1), rois(colored.size(), CV_8UC1); - rois.setTo(0); - GpuMat sub(rois, cv::Rect(rois.cols / 4, rois.rows / 4,rois.cols / 2, rois.rows / 2)); - sub.setTo(cv::Scalar::all(1)); + rois.setTo(cv::Scalar::all(1)); cv::gpu::Stream s; @@ -318,14 +304,14 @@ GPU_TEST_P(SCascadeTestAll, detectStream) cascade.detect(colored, rois, objectBoxes, s); s.waitForCompletion(); - // typedef cv::gpu::SCascade::Detection Detection; - // cv::Mat detections(objectBoxes); - // int a = *(detections.ptr(0)); - // ASSERT_EQ(a, 2448); + typedef cv::gpu::SCascade::Detection Detection; + cv::Mat detections(objectBoxes); + int a = *(detections.ptr(0)); + ASSERT_EQ(a, expected); } 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")))); + testing::Values(Fixture("cascades/inria_caltech-17.01.2013.xml", 7), + Fixture("cascades/sc_cvpr_2012_to_opencv_new_format.xml", 1291)))); #endif diff --git a/modules/objdetect/include/opencv2/objdetect/objdetect.hpp b/modules/objdetect/include/opencv2/objdetect/objdetect.hpp index 95300c48a..8d393293e 100644 --- a/modules/objdetect/include/opencv2/objdetect/objdetect.hpp +++ b/modules/objdetect/include/opencv2/objdetect/objdetect.hpp @@ -488,16 +488,6 @@ protected: Ptr maskGenerator; }; - -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. class CV_EXPORTS_W SCascade : public Algorithm {