From 8acfbde68e5b9c4bd6640a21c384c3fff8fd28aa Mon Sep 17 00:00:00 2001 From: "marina.kolpakova" Date: Wed, 14 Nov 2012 14:21:22 +0400 Subject: [PATCH] remove debug detect at scale method --- modules/gpu/doc/object_detection.rst | 4 -- modules/gpu/include/opencv2/gpu/gpu.hpp | 2 - modules/gpu/src/cuda/isf-sc.cu | 37 +++---------- modules/gpu/src/icf.hpp | 5 +- modules/gpu/src/softcascade.cpp | 35 +----------- modules/gpu/test/test_softcascade.cpp | 98 ++++++++++++++++----------------- 6 files changed, 64 insertions(+), 117 deletions(-) diff --git a/modules/gpu/doc/object_detection.rst b/modules/gpu/doc/object_detection.rst index 6434871..c503d93 100644 --- a/modules/gpu/doc/object_detection.rst +++ b/modules/gpu/doc/object_detection.rst @@ -248,7 +248,6 @@ Implementation of soft (stageless) cascaded detector. :: virtual ~SCascade(); virtual bool load(const FileNode& fn); virtual void detect(InputArray image, InputArray rois, OutputArray objects, Stream& stream = Stream::Null()) const; - virtual void detect(InputArray image, InputArray rois, OutputArray objects, const int level, Stream& stream = Stream::Null()) const; void genRoi(InputArray roi, OutputArray mask, Stream& stream = Stream::Null()) const; }; @@ -292,7 +291,6 @@ SCascade::detect Apply cascade to an input frame and return the vector of Decection objcts. .. ocv:function:: void detect(InputArray image, InputArray rois, OutputArray objects, Stream& stream = Stream::Null()) const -.. ocv:function:: void detect(InputArray image, InputArray rois, OutputArray objects, const int level, Stream& stream = Stream::Null()) const :param image: a frame on which detector will be applied. @@ -302,8 +300,6 @@ Apply cascade to an input frame and return the vector of Decection objcts. :param stream: a high-level CUDA stream abstraction used for asynchronous execution. - :param level: used for execution cascade on specific scales pyramid level. - SCascade::genRoi -------------------------- diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 9a43760..db228a6 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -1577,9 +1577,7 @@ public: // Param objects is an output array of Detections represented as GpuMat of detections (SCascade::Detection) // The first element of the matrix is actually a count of detections. // Param stream is stream is a high-level CUDA stream abstraction used for asynchronous execution - // Param level used for execution cascade on specific scales pyramid level. virtual void detect(InputArray image, InputArray rois, OutputArray objects, Stream& stream = Stream::Null()) const; - virtual void detect(InputArray image, InputArray rois, OutputArray objects, const int level, 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. diff --git a/modules/gpu/src/cuda/isf-sc.cu b/modules/gpu/src/cuda/isf-sc.cu index 7f7a10e..3d35366 100644 --- a/modules/gpu/src/cuda/isf-sc.cu +++ b/modules/gpu/src/cuda/isf-sc.cu @@ -45,15 +45,6 @@ #include #include -// #define LOG_CUDA_CASCADE - -#if defined LOG_CUDA_CASCADE -# define dprintf(format, ...) \ - do { printf(format, __VA_ARGS__); } while (0) -#else -# define dprintf(format, ...) -#endif - namespace cv { namespace gpu { namespace device { namespace icf { @@ -254,12 +245,12 @@ __global__ void soft_cascade(const CascadeInvoker invoker, Detection* ob template void CascadeInvoker::operator()(const PtrStepSzb& roi, const PtrStepSzi& hogluv, - PtrStepSz objects, PtrStepSzi counter, const int downscales, const int scale, const cudaStream_t& stream) const + PtrStepSz objects, PtrStepSzi counter, const int downscales, const cudaStream_t& stream) const { int fw = roi.rows; int fh = roi.cols; - dim3 grid(fw, fh / Policy::STA_Y, (scale == -1) ? downscales : 1); + dim3 grid(fw, fh / Policy::STA_Y, downscales); uint* ctr = (uint*)(counter.ptr(0)); Detection* det = (Detection*)objects.ptr(); @@ -268,26 +259,16 @@ void CascadeInvoker::operator()(const PtrStepSzb& roi, const PtrStepSzi& cudaChannelFormatDesc desc = cudaCreateChannelDesc(); cudaSafeCall( cudaBindTexture2D(0, thogluv, hogluv.data, desc, hogluv.cols, hogluv.rows, hogluv.step)); - cudaChannelFormatDesc desc_roi = cudaCreateChannelDesc(); - cudaSafeCall( cudaBindTexture2D(0, troi, roi.data, desc_roi, roi.cols / 8, roi.rows, roi.step)); + cudaChannelFormatDesc desc_roi = cudaCreateChannelDesc(); + cudaSafeCall( cudaBindTexture2D(0, troi, roi.data, desc_roi, roi.cols / Policy::STA_Y, roi.rows, roi.step)); const CascadeInvoker inv = *this; - if (scale == -1) - { - soft_cascade<<>>(inv, det, max_det, ctr, 0); - cudaSafeCall( cudaGetLastError()); + soft_cascade<<>>(inv, det, max_det, ctr, 0); + cudaSafeCall( cudaGetLastError()); - grid = dim3(fw, fh / Policy::STA_Y, scales - downscales); - soft_cascade<<>>(inv, det, max_det, ctr, downscales); - } - else - { - if (scale >= downscales) - soft_cascade<<>>(inv, det, max_det, ctr, scale); - else - soft_cascade<<>>(inv, det, max_det, ctr, scale); - } + grid = dim3(fw, fh / Policy::STA_Y, scales - downscales); + soft_cascade<<>>(inv, det, max_det, ctr, downscales); if (!stream) { @@ -297,7 +278,7 @@ void CascadeInvoker::operator()(const PtrStepSzb& roi, const PtrStepSzi& } template void CascadeInvoker::operator()(const PtrStepSzb& roi, const PtrStepSzi& hogluv, - PtrStepSz objects, PtrStepSzi counter, const int downscales, const int scale, const cudaStream_t& stream) const; + PtrStepSz objects, PtrStepSzi counter, const int downscales, const cudaStream_t& stream) const; } }}} \ No newline at end of file diff --git a/modules/gpu/src/icf.hpp b/modules/gpu/src/icf.hpp index 8eb080e..2bbbb64 100644 --- a/modules/gpu/src/icf.hpp +++ b/modules/gpu/src/icf.hpp @@ -118,9 +118,10 @@ struct __align__(16) Detection struct GK107PolicyX4 { enum {WARP = 32, STA_X = WARP, STA_Y = 8, SHRINKAGE = 4}; + typedef float2 roi_type; static const dim3 block() { - return dim3(GK107PolicyX4::STA_X, GK107PolicyX4::STA_Y); + return dim3(STA_X, STA_Y); } }; @@ -146,7 +147,7 @@ struct CascadeInvoker int scales; void operator()(const PtrStepSzb& roi, const PtrStepSzi& hogluv, PtrStepSz objects, - PtrStepSzi counter, const int downscales, const int csale = -1, const cudaStream_t& stream = 0) const; + PtrStepSzi counter, const int downscales, const cudaStream_t& stream = 0) const; template __device void detect(Detection* objects, const uint ndetections, uint* ctr, const int downscales) const; diff --git a/modules/gpu/src/softcascade.cpp b/modules/gpu/src/softcascade.cpp index ad6e000..5da3abf 100644 --- a/modules/gpu/src/softcascade.cpp +++ b/modules/gpu/src/softcascade.cpp @@ -311,13 +311,13 @@ struct cv::gpu::SCascade::Fields leaves.upload(hleaves); } - void detect(int scale, const cv::gpu::GpuMat& roi, const cv::gpu::GpuMat& count, cv::gpu::GpuMat& objects, const cudaStream_t& stream) const + void detect(const cv::gpu::GpuMat& roi, const cv::gpu::GpuMat& count, cv::gpu::GpuMat& objects, const cudaStream_t& stream) const { cudaMemset(count.data, 0, sizeof(Detection)); cudaSafeCall( cudaGetLastError()); device::icf::CascadeInvoker invoker = device::icf::CascadeInvoker(levels, octaves, stages, nodes, leaves); - invoker(roi, hogluv, objects, count, downscales, scale, stream); + invoker(roi, hogluv, objects, count, downscales, stream); } void preprocess(const cv::gpu::GpuMat& colored, Stream& s) @@ -521,36 +521,7 @@ void cv::gpu::SCascade::detect(InputArray image, InputArray _rois, OutputArray _ objects = GpuMat(objects, cv::Rect( sizeof(Detection), 0, objects.cols - sizeof(Detection), 1)); cudaStream_t stream = StreamAccessor::getStream(s); - flds.detect(-1, rois, tmp, objects, stream); -} - -void cv::gpu::SCascade::detect(InputArray image, InputArray _rois, OutputArray _objects, const int level, Stream& s) const -{ - CV_Assert(fields); - - const GpuMat colored = image.getGpuMat(); - // only color images are supperted - CV_Assert(colored.type() == CV_8UC3 || colored.type() == CV_32SC1); - - Fields& flds = *fields; - if (colored.type() == CV_8UC3) - { - // only this window size allowed - // CV_Assert(colored.cols == Fields::FRAME_WIDTH && colored.rows == Fields::FRAME_HEIGHT); - flds.preprocess(colored, s); - } - else - { - colored.copyTo(flds.hogluv); - } - - GpuMat rois = _rois.getGpuMat(), objects = _objects.getGpuMat(); - - GpuMat tmp = GpuMat(objects, cv::Rect(0, 0, sizeof(Detection), 1)); - objects = GpuMat(objects, cv::Rect( sizeof(Detection), 0, objects.cols - sizeof(Detection), 1)); - cudaStream_t stream = StreamAccessor::getStream(s); - - flds.detect(level, rois, tmp, objects, stream); + flds.detect(rois, tmp, objects, stream); } void cv::gpu::SCascade::genRoi(InputArray _roi, OutputArray _mask, Stream& stream) const diff --git a/modules/gpu/test/test_softcascade.cpp b/modules/gpu/test/test_softcascade.cpp index 7034b33..e36c289 100644 --- a/modules/gpu/test/test_softcascade.cpp +++ b/modules/gpu/test/test_softcascade.cpp @@ -195,55 +195,55 @@ GPU_TEST_P(SCascadeTestRoi, detect, } -typedef ::testing::TestWithParam > SCascadeTestLevel; -GPU_TEST_P(SCascadeTestLevel, detect, - 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::Range(0, 47) - )) -{ - cv::gpu::setDevice(GET_PARAM(0).deviceID()); - - cv::gpu::SCascade cascade; - - cv::FileStorage fs(perf::TestBase::getDataPath(GET_PARAM(1)), cv::FileStorage::READ); - ASSERT_TRUE(fs.isOpened()); - - ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); - - cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() + GET_PARAM(2)); - ASSERT_FALSE(coloredCpu.empty()); - - typedef cv::gpu::SCascade::Detection Detection; - GpuMat colored(coloredCpu), objectBoxes(1, 100 * sizeof(Detection), CV_8UC1), rois(colored.size(), CV_8UC1); - rois.setTo(1); - - cv::gpu::GpuMat trois; - cascade.genRoi(rois, trois); - objectBoxes.setTo(0); - int level = GET_PARAM(3); - cascade.detect(colored, trois, objectBoxes, level); - - cv::Mat dt(objectBoxes); - - Detection* dts = ((Detection*)dt.data) + 1; - int* count = dt.ptr(0); - - cv::Mat result(coloredCpu); - - printTotal(std::cout, *count); - for (int i = 0; i < *count; ++i) - { - Detection d = dts[i]; - print(std::cout, d); - cv::rectangle(result, cv::Rect(d.x, d.y, d.w, d.h), cv::Scalar(255, 0, 0, 255), 1); - } - - writeResult(result, level); - SHOW(result); -} +// typedef ::testing::TestWithParam > SCascadeTestLevel; +// GPU_TEST_P(SCascadeTestLevel, detect, +// 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::Range(0, 47) +// )) +// { +// cv::gpu::setDevice(GET_PARAM(0).deviceID()); + +// cv::gpu::SCascade cascade; + +// cv::FileStorage fs(perf::TestBase::getDataPath(GET_PARAM(1)), cv::FileStorage::READ); +// ASSERT_TRUE(fs.isOpened()); + +// ASSERT_TRUE(cascade.load(fs.getFirstTopLevelNode())); + +// cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() + GET_PARAM(2)); +// ASSERT_FALSE(coloredCpu.empty()); + +// typedef cv::gpu::SCascade::Detection Detection; +// GpuMat colored(coloredCpu), objectBoxes(1, 100 * sizeof(Detection), CV_8UC1), rois(colored.size(), CV_8UC1); +// rois.setTo(1); + +// cv::gpu::GpuMat trois; +// cascade.genRoi(rois, trois); +// objectBoxes.setTo(0); +// int level = GET_PARAM(3); +// cascade.detect(colored, trois, objectBoxes, level); + +// cv::Mat dt(objectBoxes); + +// Detection* dts = ((Detection*)dt.data) + 1; +// int* count = dt.ptr(0); + +// cv::Mat result(coloredCpu); + +// printTotal(std::cout, *count); +// for (int i = 0; i < *count; ++i) +// { +// Detection d = dts[i]; +// print(std::cout, d); +// cv::rectangle(result, cv::Rect(d.x, d.y, d.w, d.h), cv::Scalar(255, 0, 0, 255), 1); +// } + +// writeResult(result, level); +// SHOW(result); +// } TEST(SCascadeTest, readCascade) { -- 2.7.4