LBP classifier: tracking of count of detected objects was moved in cascadeclassifier.cpp

diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp
index f3b3884..2057f1a 100644 (file)
--- a/modules/gpu/src/cascadeclassifier.cpp
+++ b/modules/gpu/src/cascadeclassifier.cpp
@@ -273,21 +273,22 @@ namespace cv { namespace gpu { namespace device
 {\r
     namespace lbp\r
     {\r
-        int classifyStump(const DevMem2Db mstages,\r
-                           const int nstages,\r
-                           const DevMem2Di mnodes,\r
-                           const DevMem2Df mleaves,\r
-                           const DevMem2Di msubsets,\r
-                           const DevMem2Db mfeatures,\r
-                           const DevMem2Di integral,\r
-                           const int workWidth,\r
-                           const int workHeight,\r
-                           const int clWidth,\r
-                           const int clHeight,\r
-                           float scale,\r
-                           int step,\r
-                           int subsetSize,\r
-                           DevMem2D_<int4> objects);\r
+        classifyStump(const DevMem2Db mstages,\r
+                      const int nstages,\r
+                      const DevMem2Di mnodes,\r
+                      const DevMem2Df mleaves,\r
+                      const DevMem2Di msubsets,\r
+                      const DevMem2Db mfeatures,\r
+                      const DevMem2Di integral,\r
+                      const int workWidth,\r
+                      const int workHeight,\r
+                      const int clWidth,\r
+                      const int clHeight,\r
+                      float scale,\r
+                      int step,\r
+                      int subsetSize,\r
+                      DevMem2D_<int4> objects,\r
+                      unsigned int* classified);\r
     }\r
 }}}\r
 \r
@@ -308,6 +309,11 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp
         maxObjectSize = image.size();\r
 \r
     scaledImageBuffer.create(image.rows + 1, image.cols + 1, CV_8U);\r
+    unsigned int* classified = new unsigned int[1];\r
+    *classified = 0;\r
+    unsigned int* dclassified;\r
+    cudaMalloc(&dclassified, sizeof(int));\r
+    cudaMemcpy(dclassified, classified, sizeof(int), cudaMemcpyHostToDevice);\r
 \r
     for( double factor = 1; ; factor *= scaleFactor )\r
     {\r
@@ -331,10 +337,11 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp
 \r
         int step = (factor <= 2.) + 1;\r
 \r
-        int res = cv::gpu::device::lbp::classifyStump(stage_mat, stage_mat.cols / sizeof(Stage), nodes_mat, leaves_mat, subsets_mat, features_mat,\r
-        integral, processingRectSize.width, processingRectSize.height, windowSize.width, windowSize.height, scaleFactor, step, subsetSize, objects);\r
-        std::cout  << res << "Results:    " << cv::Mat(objects).row(0).colRange(0, res) << std::endl;\r
+        cv::gpu::device::lbp::classifyStump(stage_mat, stage_mat.cols / sizeof(Stage), nodes_mat, leaves_mat, subsets_mat, features_mat,\r
+        integral, processingRectSize.width, processingRectSize.height, windowSize.width, windowSize.height, scaleFactor, step, subsetSize, objects, dclassified);\r
     }\r
+        cudaMemcpy(classified, dclassified, sizeof(int), cudaMemcpyDeviceToHost);\r
+        std::cout  << *classified << "Results:    " << cv::Mat(objects).row(0).colRange(0, *classified) << std::endl;\r
     // TODO: reject levels\r
 \r
     return 0;\r

diff --git a/modules/gpu/src/cuda/lbp.cu b/modules/gpu/src/cuda/lbp.cu
index b997947..b07ecad 100644 (file)
--- a/modules/gpu/src/cuda/lbp.cu
+++ b/modules/gpu/src/cuda/lbp.cu
@@ -51,8 +51,6 @@ namespace cv { namespace gpu { namespace device
         {
             int y = threadIdx.x * scale;
             int x = blockIdx.x * scale;
-            *n = 0;
-            int i = 0;
 
             int current_node = 0;
             int current_leave = 0;
@@ -77,7 +75,6 @@ namespace cv { namespace gpu { namespace device
                     current_leave += 2;
                 }
 
-                i = s;
                 if (sum < stage.threshold)
                     return;
             }
@@ -88,29 +85,26 @@ namespace cv { namespace gpu { namespace device
             rect.z = roundf(clWidth);
             rect.w = roundf(clHeight);
 
-            int res = atomicInc(n, 1000);
+            int res = atomicInc(n, 100);
             objects(0, res) = rect;
         }
 
-        int classifyStump(const DevMem2Db mstages, const int nstages, const DevMem2Di mnodes, const DevMem2Df mleaves, const DevMem2Di msubsets, const DevMem2Db mfeatures,
+        classifyStump(const DevMem2Db mstages, const int nstages, const DevMem2Di mnodes, const DevMem2Df mleaves, const DevMem2Di msubsets, const DevMem2Db mfeatures,
                            const DevMem2Di integral, const int workWidth, const int workHeight, const int clWidth, const int clHeight, float scale, int step, int subsetSize,
-                           DevMem2D_<int4> objects)
+                           DevMem2D_<int4> objects, unsigned int* classified)
         {
             int blocks  = ceilf(workHeight / (float)step);
             int threads = ceilf(workWidth / (float)step);
-            printf("blocks %d, threads %d\n", blocks, threads);
+            // printf("blocks %d, threads %d\n", blocks, threads);
 
             Stage* stages = (Stage*)(mstages.ptr());
             ClNode* nodes = (ClNode*)(mnodes.ptr());
             const float* leaves = mleaves.ptr();
             const int* subsets = msubsets.ptr();
             const uchar4* features = (uchar4*)(mfeatures.ptr());
-            unsigned int * n, *h_n = new unsigned int[1];
-            cudaMalloc(&n, sizeof(int));
+
             lbp_classify_stump<<<blocks, threads>>>(stages, nstages, nodes, leaves, subsets, features, integral,
-                workWidth, workHeight, clWidth, clHeight, scale, step, subsetSize, objects, n);
-            cudaMemcpy(h_n, n, sizeof(int), cudaMemcpyDeviceToHost);
-            return *h_n;
+                workWidth, workHeight, clWidth, clHeight, scale, step, subsetSize, objects, classified);
         }
     }
 }}}
\ No newline at end of file