add roi support

diff --git a/modules/gpu/perf/perf_objdetect.cpp b/modules/gpu/perf/perf_objdetect.cpp
index 8531372..a863371 100644 (file)
--- a/modules/gpu/perf/perf_objdetect.cpp
+++ b/modules/gpu/perf/perf_objdetect.cpp
@@ -104,7 +104,11 @@ PERF_TEST_P(SoftCascade, detect, Values<pair_string>(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 (file)
--- a/modules/gpu/src/cuda/isf-sc.cu
+++ b/modules/gpu/src/cuda/isf-sc.cu
@@ -86,6 +86,7 @@ namespace icf {
     }
 
     texture<int,  cudaTextureType2D, cudaReadModeElementType> thogluv;
+    texture<char,  cudaTextureType2D, cudaReadModeElementType> troi;
 
     template<bool isUp>
     __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<uchar4> objects, PtrStepSzi counter, const int downscales)
     {
@@ -350,6 +357,9 @@ namespace icf {
         cudaChannelFormatDesc desc = cudaCreateChannelDesc<int>();
         cudaSafeCall( cudaBindTexture2D(0, thogluv, hogluv.data, desc, hogluv.cols, hogluv.rows, hogluv.step));
 
+        cudaChannelFormatDesc desc_roi = cudaCreateChannelDesc<char>();
+        cudaSafeCall( cudaBindTexture2D(0, troi, roi.data, desc_roi, roi.cols, roi.rows, roi.step));
+
         test_kernel_warp<false><<<grid, block>>>(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<uchar4> 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<uchar4> objects, PtrStepSzi counter, const int downscales)
     {
         int fw = 160;
         int fh = 120;
@@ -381,6 +391,9 @@ namespace icf {
         cudaChannelFormatDesc desc = cudaCreateChannelDesc<int>();
         cudaSafeCall( cudaBindTexture2D(0, thogluv, hogluv.data, desc, hogluv.cols, hogluv.rows, hogluv.step));
 
+        cudaChannelFormatDesc desc_roi = cudaCreateChannelDesc<char>();
+        cudaSafeCall( cudaBindTexture2D(0, troi, roi.data, desc_roi, roi.cols, roi.rows, roi.step));
+
         if (scale >= downscales)
             test_kernel_warp<true><<<grid, block>>>(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 (file)
--- 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<uchar4> 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<uchar4> 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<uchar4> 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<uchar4> 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 (file)
--- 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<std::tr1::tuple<std::string, int> > {
@@ -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);