From 319c20c797704b33bae284af9f5f11fd0da99697 Mon Sep 17 00:00:00 2001
From: Marina Kolpakova <no@email>
Date: Mon, 2 Jul 2012 08:08:17 +0000
Subject: [PATCH] result storing: atomic based

---
 modules/gpu/src/cascadeclassifier.cpp | 16 +++++++---------
 modules/gpu/src/cuda/lbp.cu           | 18 ++++++++++--------
 2 files changed, 17 insertions(+), 17 deletions(-)

diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp
index e1a4554ff..556d92d50 100644
--- a/modules/gpu/src/cascadeclassifier.cpp
+++ b/modules/gpu/src/cascadeclassifier.cpp
@@ -272,7 +272,7 @@ namespace cv { namespace gpu { namespace device
 {
     namespace lbp
     {
-        void classifyStump(const DevMem2Db mstages,
+        int classifyStump(const DevMem2Db mstages,
                            const int nstages,
                            const DevMem2Di mnodes,
                            const DevMem2Df mleaves,
@@ -298,13 +298,10 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp
 
     const int defaultObjSearchNum = 100;
 
-    // if( !objects.empty() && objects.depth() == CV_32S)
-    //     objects.reshape(4, 1);
-    // else
-    //     objects.create(1 , defaultObjSearchNum, CV_32SC4);
-
-    // temp solution
-    objects.create(image.rows, image.cols, CV_32SC4);
+    if( !objects.empty() && objects.depth() == CV_32S)
+        objects.reshape(4, 1);
+    else
+        objects.create(1 , defaultObjSearchNum, CV_32SC4);
 
     if (maxObjectSize == cv::Size())
         maxObjectSize = image.size();
@@ -333,8 +330,9 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp
 
         int step = (factor <= 2.) + 1;
 
-        cv::gpu::device::lbp::classifyStump(stage_mat, stage_mat.cols / sizeof(Stage), nodes_mat, leaves_mat, subsets_mat, features_mat,
+        int res = cv::gpu::device::lbp::classifyStump(stage_mat, stage_mat.cols / sizeof(Stage), nodes_mat, leaves_mat, subsets_mat, features_mat,
         integral, processingRectSize.width, processingRectSize.height, windowSize.width, windowSize.height, scaleFactor, step, subsetSize, objects);
+        std::cout  << res << "Results:    " << cv::Mat(objects).row(0).colRange(0, res) << std::endl;
     }
     // TODO: reject levels
 
diff --git a/modules/gpu/src/cuda/lbp.cu b/modules/gpu/src/cuda/lbp.cu
index 713392045..b9979474f 100644
--- a/modules/gpu/src/cuda/lbp.cu
+++ b/modules/gpu/src/cuda/lbp.cu
@@ -47,11 +47,11 @@ namespace cv { namespace gpu { namespace device
     namespace lbp
     {
         __global__ void lbp_classify_stump(Stage* stages, int nstages, ClNode* nodes, const float* leaves, const int* subsets, const uchar4* features,
-            const DevMem2Di integral, int workWidth, int workHeight, int clWidth, int clHeight, float scale, int step, int subsetSize, DevMem2D_<int4> objects)
+            const DevMem2Di integral, int workWidth, int workHeight, int clWidth, int clHeight, float scale, int step, int subsetSize, DevMem2D_<int4> objects, unsigned int* n)
         {
             int y = threadIdx.x * scale;
             int x = blockIdx.x * scale;
-
+            *n = 0;
             int i = 0;
 
             int current_node = 0;
@@ -88,12 +88,11 @@ namespace cv { namespace gpu { namespace device
             rect.z = roundf(clWidth);
             rect.w = roundf(clHeight);
 
-            if(i >= 19)
-                printf( "GPU detected [%d, %d] - [%d, %d]\n", rect.x, rect.y, rect.z, rect.w);
-
+            int res = atomicInc(n, 1000);
+            objects(0, res) = rect;
         }
 
-        void classifyStump(const DevMem2Db mstages, const int nstages, const DevMem2Di mnodes, const DevMem2Df mleaves, const DevMem2Di msubsets, const DevMem2Db mfeatures,
+        int 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)
         {
@@ -106,9 +105,12 @@ namespace cv { namespace gpu { namespace device
             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);
+                workWidth, workHeight, clWidth, clHeight, scale, step, subsetSize, objects, n);
+            cudaMemcpy(h_n, n, sizeof(int), cudaMemcpyDeviceToHost);
+            return *h_n;
         }
     }
 }}}
\ No newline at end of file