From: marina.kolpakova Date: Tue, 2 Oct 2012 13:25:26 +0000 (+0400) Subject: add roi support X-Git-Tag: submit/tizen_ivi/20141117.190038~2^2~1214^2~59^2~42 X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=eb91593c08daca478663d7d0805f0ce664fd5d9b;p=profile%2Fivi%2Fopencv.git add roi support --- diff --git a/modules/gpu/perf/perf_objdetect.cpp b/modules/gpu/perf/perf_objdetect.cpp index 8531372..a863371 100644 --- a/modules/gpu/perf/perf_objdetect.cpp +++ b/modules/gpu/perf/perf_objdetect.cpp @@ -104,7 +104,11 @@ PERF_TEST_P(SoftCascade, detect, Values(make_pair("cv/cascadeandhog cv::gpu::SoftCascade cascade; ASSERT_TRUE(cascade.load(perf::TestBase::getDataPath(GetParam().first))); - cv::gpu::GpuMat rois, objectBoxes(1, 16384, CV_8UC1); + cv::gpu::GpuMat objectBoxes(1, 16384, CV_8UC1), rois(cascade.getRoiSize(), CV_8UC1); + + rois.setTo(0); + cv::gpu::GpuMat sub(rois, cv::Rect(rois.cols / 4, rois.rows / 4,rois.cols / 2, rois.rows / 2)); + sub.setTo(cv::Scalar::all(1)); cascade.detectMultiScale(colored, rois, objectBoxes); TEST_CYCLE() diff --git a/modules/gpu/src/cuda/isf-sc.cu b/modules/gpu/src/cuda/isf-sc.cu index c9a92e3..4bde7f7 100644 --- a/modules/gpu/src/cuda/isf-sc.cu +++ b/modules/gpu/src/cuda/isf-sc.cu @@ -86,6 +86,7 @@ namespace icf { } texture thogluv; + texture troi; template __device__ __forceinline__ float rescale(const Level& level, Node& node) @@ -213,6 +214,8 @@ namespace icf { if(x >= level.workRect.x || y >= level.workRect.y) return; + if (!tex2D(troi, x, y)) return; + Octave octave = octaves[level.octave]; int st = octave.index * octave.stages; const int stEnd = st + 1024; @@ -279,6 +282,10 @@ namespace icf { // if (blockIdx.z != 31) return; if(x >= level.workRect.x || y >= level.workRect.y) return; + int roi = tex2D(troi, x, y); + printf("%d\n", roi); + if (!roi) return; + Octave octave = octaves[level.octave]; int st = octave.index * octave.stages; @@ -328,7 +335,7 @@ namespace icf { } #endif - void detect(const PtrStepSzb& levels, const PtrStepSzb& octaves, const PtrStepSzf& stages, + void detect(const PtrStepSzb& roi, const PtrStepSzb& levels, const PtrStepSzb& octaves, const PtrStepSzf& stages, const PtrStepSzb& nodes, const PtrStepSzf& leaves, const PtrStepSzi& hogluv, PtrStepSz objects, PtrStepSzi counter, const int downscales) { @@ -350,6 +357,9 @@ namespace icf { 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, roi.rows, roi.step)); + test_kernel_warp<<>>(l, oct, st, nd, lf, det, max_det, ctr, 0); cudaSafeCall( cudaGetLastError()); @@ -359,9 +369,9 @@ namespace icf { cudaSafeCall( cudaDeviceSynchronize()); } - void detectAtScale(const int scale, const PtrStepSzb& levels, const PtrStepSzb& octaves, const PtrStepSzf& stages, - const PtrStepSzb& nodes, const PtrStepSzf& leaves, const PtrStepSzi& hogluv, PtrStepSz objects, - PtrStepSzi counter, const int downscales) + void detectAtScale(const int scale, const PtrStepSzb& roi, const PtrStepSzb& levels, const PtrStepSzb& octaves, + const PtrStepSzf& stages, const PtrStepSzb& nodes, const PtrStepSzf& leaves, const PtrStepSzi& hogluv, + PtrStepSz objects, PtrStepSzi counter, const int downscales) { int fw = 160; int fh = 120; @@ -381,6 +391,9 @@ namespace icf { 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, roi.rows, roi.step)); + if (scale >= downscales) test_kernel_warp<<>>(l, oct, st, nd, lf, det, max_det, ctr, scale); else diff --git a/modules/gpu/src/softcascade.cpp b/modules/gpu/src/softcascade.cpp index af83669..9ea365c 100644 --- a/modules/gpu/src/softcascade.cpp +++ b/modules/gpu/src/softcascade.cpp @@ -69,12 +69,29 @@ namespace cv { namespace gpu { namespace device { namespace icf { void fillBins(cv::gpu::PtrStepSzb hogluv, const cv::gpu::PtrStepSzf& nangle, const int fw, const int fh, const int bins); - void detect(const PtrStepSzb& levels, const PtrStepSzb& octaves, const PtrStepSzf& stages, - const PtrStepSzb& nodes, const PtrStepSzf& leaves, const PtrStepSzi& hogluv, PtrStepSz objects, - PtrStepSzi counter, const int downscales); - void detectAtScale(const int scale, const PtrStepSzb& levels, const PtrStepSzb& octaves, const PtrStepSzf& stages, - const PtrStepSzb& nodes, const PtrStepSzf& leaves, const PtrStepSzi& hogluv, PtrStepSz objects, - PtrStepSzi counter, const int downscales); + + void detect(const PtrStepSzb& rois, + const PtrStepSzb& levels, + const PtrStepSzb& octaves, + const PtrStepSzf& stages, + const PtrStepSzb& nodes, + const PtrStepSzf& leaves, + const PtrStepSzi& hogluv, + PtrStepSz objects, + PtrStepSzi counter, + const int downscales); + + void detectAtScale(const int scale, + const PtrStepSzb& rois, + const PtrStepSzb& levels, + const PtrStepSzb& octaves, + const PtrStepSzf& stages, + const PtrStepSzb& nodes, + const PtrStepSzf& leaves, + const PtrStepSzi& hogluv, + PtrStepSz objects, + PtrStepSzi counter, + const int downscales); } }}} @@ -143,16 +160,16 @@ struct cv::gpu::SoftCascade::Filds }; bool fill(const FileNode &root, const float mins, const float maxs); - void detect(cv::gpu::GpuMat objects, cudaStream_t stream) const + void detect(cv::gpu::GpuMat roi, cv::gpu::GpuMat objects, cudaStream_t stream) const { cudaMemset(detCounter.data, 0, detCounter.step * detCounter.rows * sizeof(int)); - device::icf::detect(levels, octaves, stages, nodes, leaves, hogluv, objects , detCounter, downscales); + device::icf::detect(roi, levels, octaves, stages, nodes, leaves, hogluv, objects , detCounter, downscales); } - void detectAtScale(int scale, cv::gpu::GpuMat objects, cudaStream_t stream) const + void detectAtScale(int scale, cv::gpu::GpuMat roi, cv::gpu::GpuMat objects, cudaStream_t stream) const { cudaMemset(detCounter.data, 0, detCounter.step * detCounter.rows * sizeof(int)); - device::icf::detectAtScale(scale, levels, octaves, stages, nodes, leaves, hogluv, objects, + device::icf::detectAtScale(scale, roi, levels, octaves, stages, nodes, leaves, hogluv, objects, detCounter, downscales); } @@ -467,6 +484,9 @@ void cv::gpu::SoftCascade::detectMultiScale(const GpuMat& colored, const GpuMat& // only color images are supperted CV_Assert(colored.type() == CV_8UC3); + // we guess user knows about shrincage + CV_Assert((rois.size() == getRoiSize()) && (rois.type() == CV_8UC1)); + // only this window size allowed CV_Assert(colored.cols == Filds::FRAME_WIDTH && colored.rows == Filds::FRAME_HEIGHT); @@ -551,9 +571,9 @@ void cv::gpu::SoftCascade::detectMultiScale(const GpuMat& colored, const GpuMat& #endif if (specificScale == -1) - flds.detect(objects, 0); + flds.detect(rois,objects, 0); else - flds.detectAtScale(specificScale, objects, 0); + flds.detectAtScale(specificScale, rois, objects, 0); cv::Mat out(flds.detCounter); int ndetections = *(out.data); diff --git a/modules/gpu/test/test_softcascade.cpp b/modules/gpu/test/test_softcascade.cpp index 4d1a4b7..8460291 100644 --- a/modules/gpu/test/test_softcascade.cpp +++ b/modules/gpu/test/test_softcascade.cpp @@ -63,12 +63,13 @@ TEST(SoftCascade, detect) cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() + "../cv/cascadeandhog/bahnhof/image_00000000_0.png"); ASSERT_FALSE(coloredCpu.empty()); - GpuMat colored(coloredCpu), objectBoxes(1, 100000, CV_8UC1), rois; - // ASSERT_NO_THROW( - // { - cascade.detectMultiScale(colored, rois, objectBoxes); - // }); + GpuMat colored(coloredCpu), objectBoxes(1, 100000, CV_8UC1), rois(cascade.getRoiSize(), 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)); + + cascade.detectMultiScale(colored, rois, objectBoxes); } class SCSpecific : public ::testing::TestWithParam > { @@ -93,7 +94,10 @@ TEST_P(SCSpecific, detect) cv::Mat coloredCpu = cv::imread(cvtest::TS::ptr()->get_data_path() + path); ASSERT_FALSE(coloredCpu.empty()); - GpuMat colored(coloredCpu), objectBoxes(1, 1000, CV_8UC1), rois; + GpuMat colored(coloredCpu), objectBoxes(1, 1000, CV_8UC1), rois(cascade.getRoiSize(), 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)); int level = GET_PARAM(1); cascade.detectMultiScale(colored, rois, objectBoxes, 1, level);