added lbp cascade test, fixed race conditions problems

diff --git a/modules/gpu/src/cascadeclassifier.cpp b/modules/gpu/src/cascadeclassifier.cpp
index d2f4120..da0bfe5 100644 (file)
--- a/modules/gpu/src/cascadeclassifier.cpp
+++ b/modules/gpu/src/cascadeclassifier.cpp
@@ -290,7 +290,7 @@ namespace cv { namespace gpu { namespace device
                       DevMem2D_<int4> objects,\r
                       unsigned int* classified);\r
 \r
-        int connectedConmonents(DevMem2D_<int4> candidates, int groupThreshold, float grouping_eps, unsigned int* nclasses);\r
+        int connectedConmonents(DevMem2D_<int4> candidates, DevMem2D_<int4> objects,int groupThreshold, float grouping_eps, unsigned int* nclasses);\r
     }\r
 }}}\r
 \r
@@ -308,6 +308,7 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp
     else\r
         objects.create(1 , defaultObjSearchNum, CV_32SC4);\r
 \r
+    GpuMat candidates(1 , defaultObjSearchNum, CV_32SC4);\r
     if (maxObjectSize == cv::Size())\r
         maxObjectSize = image.size();\r
 \r
@@ -317,6 +318,7 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp
     unsigned int* dclassified;\r
     cudaMalloc(&dclassified, sizeof(int));\r
     cudaMemcpy(dclassified, classified, sizeof(int), cudaMemcpyHostToDevice);\r
+    int step;\r
 \r
     for( double factor = 1; ; factor *= scaleFactor )\r
     {\r
@@ -334,25 +336,22 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp
         //     continue;\r
 \r
         cv::gpu::resize(image, scaledImageBuffer, scaledImageSize, 0, 0, CV_INTER_LINEAR);\r
-\r
-        integral.create(cv::Size(scaledImageSize.width + 1, scaledImageSize.height + 1), CV_32SC1);\r
         cv::gpu::integral(scaledImageBuffer, integral);\r
 \r
-        int step = (factor <= 2.) + 1;\r
+        step = (factor <= 2.) + 1;\r
 \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
+        integral, processingRectSize.width, processingRectSize.height, windowSize.width, windowSize.height, factor, step, subsetSize, candidates, dclassified);\r
     }\r
-\r
-    cudaMemcpy(classified, dclassified, sizeof(int), cudaMemcpyDeviceToHost);\r
-    GpuMat candidates(1, *classified, objects.type(), objects.ptr());\r
-    // std::cout  << *classified << " Results: " << cv::Mat(candidates) << std::endl;\r
-\r
     if (groupThreshold <= 0  || objects.empty())\r
         return 0;\r
-    cv::gpu::device::lbp::connectedConmonents(candidates, groupThreshold, grouping_eps, dclassified);\r
+    cv::gpu::device::lbp::connectedConmonents(candidates, objects, groupThreshold, grouping_eps, dclassified);\r
+    cudaMemcpy(classified, dclassified, sizeof(int), cudaMemcpyDeviceToHost);\r
     cudaSafeCall( cudaDeviceSynchronize() );\r
-    return *classified;\r
+    step = *classified;\r
+    delete[] classified;\r
+    cudaFree(dclassified);\r
+    return step;\r
 }\r
 \r
 // ============ old fashioned haar cascade ==============================================//\r

diff --git a/modules/gpu/src/cuda/lbp.cu b/modules/gpu/src/cuda/lbp.cu
index 5c273b3..cd46945 100644 (file)
--- a/modules/gpu/src/cuda/lbp.cu
+++ b/modules/gpu/src/cuda/lbp.cu
@@ -51,8 +51,8 @@ namespace cv { namespace gpu { namespace device
         __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, unsigned int* n)
         {
-            int y = threadIdx.x * scale;
-            int x = blockIdx.x * scale;
+            int x = threadIdx.x * step;
+            int y = blockIdx.x * step;
 
             int current_node = 0;
             int current_leave = 0;
@@ -92,7 +92,7 @@ namespace cv { namespace gpu { namespace device
         }
 
         template<typename Pr>
-        __global__ void disjoin(int4* candidates, unsigned int n, int groupThreshold, float grouping_eps, unsigned int* nclasses)
+        __global__ void disjoin(int4* candidates, int4* objects, unsigned int n, int groupThreshold, float grouping_eps, unsigned int* nclasses)
         {
             using cv::gpu::device::VecTraits;
             unsigned int tid = threadIdx.x;
@@ -119,7 +119,7 @@ namespace cv { namespace gpu { namespace device
             __syncthreads();
 
             atomicInc((unsigned int*)labels + cls, n);
-            labels[n - 1] = 0;
+            *nclasses = 0;
 
             int active = labels[tid];
             if (active)
@@ -152,11 +152,9 @@ namespace cv { namespace gpu { namespace device
                             (n2 > max(3, n1) || n1 < 3) )
                             break;
                     }
-
                     if( j == n)
                     {
-                        // printf("founded gpu %d %d %d %d \n", r1[0], r1[1], r1[2], r1[3]);
-                        candidates[atomicInc((unsigned int*)labels + n -1, n)] = VecTraits<int4>::make(r1[0], r1[1], r1[2], r1[3]);
+                        objects[atomicInc(nclasses, n)] = VecTraits<int4>::make(r1[0], r1[1], r1[2], r1[3]);
                     }
                 }
             }
@@ -179,11 +177,11 @@ namespace cv { namespace gpu { namespace device
                 workWidth, workHeight, clWidth, clHeight, scale, step, subsetSize, objects, classified);
         }
 
-        int connectedConmonents(DevMem2D_<int4> candidates, int groupThreshold, float grouping_eps, unsigned int* nclasses)
+        int connectedConmonents(DevMem2D_<int4> candidates, DevMem2D_<int4> objects, int groupThreshold, float grouping_eps, unsigned int* nclasses)
         {
             int threads = candidates.cols;
             int smem_amount = threads * sizeof(int) + threads * sizeof(int4);
-            disjoin<InSameComponint><<<1, threads, smem_amount>>>((int4*)candidates.ptr(), candidates.cols, groupThreshold, grouping_eps, nclasses);
+            disjoin<InSameComponint><<<1, threads, smem_amount>>>((int4*)candidates.ptr(), (int4*)objects.ptr(), candidates.cols, groupThreshold, grouping_eps, nclasses);
             return 0;
         }
     }

diff --git a/modules/gpu/src/opencv2/gpu/device/lbp.hpp b/modules/gpu/src/opencv2/gpu/device/lbp.hpp
index 2b620b5..f4ec78b 100644 (file)
--- a/modules/gpu/src/opencv2/gpu/device/lbp.hpp
+++ b/modules/gpu/src/opencv2/gpu/device/lbp.hpp
@@ -65,12 +65,12 @@ namespace lbp{
     struct InSameComponint
     {
     public:
-        __device__ __forceinline__ InSameComponint(float _eps) : eps(_eps * 0.5) {}
+        __device__ __forceinline__ InSameComponint(float _eps) : eps(_eps) {}
         __device__ __forceinline__ InSameComponint(const InSameComponint& other) : eps(other.eps) {}
 
         __device__ __forceinline__ bool operator()(const int4& r1, const int4& r2) const
         {
-            double delta = eps * (min(r1.z, r2.z) + min(r1.w, r2.w));
+            float delta = eps * (min(r1.z, r2.z) + min(r1.w, r2.w)) * 0.5;
 
             return abs(r1.x - r2.x) <= delta && abs(r1.y - r2.y) <= delta
                 && abs(r1.x + r1.z - r2.x - r2.z) <= delta && abs(r1.y + r1.w - r2.y - r2.w) <= delta;

diff --git a/modules/gpu/test/test_objdetect.cpp b/modules/gpu/test/test_objdetect.cpp
index 9c6db09..8b49538 100644 (file)
--- a/modules/gpu/test/test_objdetect.cpp
+++ b/modules/gpu/test/test_objdetect.cpp
@@ -308,4 +308,57 @@ INSTANTIATE_TEST_CASE_P(GPU_ObjDetect, LBP_Read_classifier, testing::Combine(
     testing::Values<int>(0)\r
     ));\r
 \r
+PARAM_TEST_CASE(LBP_classify, cv::gpu::DeviceInfo, int)\r
+{\r
+    cv::gpu::DeviceInfo devInfo;\r
+\r
+    virtual void SetUp()\r
+    {\r
+        devInfo = GET_PARAM(0);\r
+        cv::gpu::setDevice(devInfo.deviceID());\r
+    }\r
+};\r
+\r
+TEST_P(LBP_classify, Accuracy)\r
+{\r
+    std::string classifierXmlPath = std::string(cvtest::TS::ptr()->get_data_path()) + "lbpcascade/lbpcascade_frontalface.xml";\r
+    std::string imagePath = std::string(cvtest::TS::ptr()->get_data_path()) + "lbpcascade/er.png";\r
+\r
+    cv::CascadeClassifier cpuClassifier(classifierXmlPath);\r
+    ASSERT_FALSE(cpuClassifier.empty());\r
+\r
+    cv::Mat image = cv::imread(imagePath);\r
+    image = image.colRange(0, image.cols / 2);\r
+    cv::Mat grey;\r
+    cvtColor(image, grey, CV_BGR2GRAY);\r
+    ASSERT_FALSE(image.empty());\r
+\r
+    std::vector<cv::Rect> rects;\r
+    cpuClassifier.detectMultiScale(grey, rects);\r
+    cv::Mat markedImage = image.clone();\r
+\r
+    std::vector<cv::Rect>::iterator it = rects.begin();\r
+    for (; it != rects.end(); ++it)\r
+        cv::rectangle(markedImage, *it, cv::Scalar(255, 0, 0, 255));\r
+\r
+    cv::gpu::CascadeClassifier_GPU_LBP gpuClassifier;\r
+    ASSERT_TRUE(gpuClassifier.load(classifierXmlPath));\r
+    cv::gpu::GpuMat gpu_rects, buffer;\r
+    cv::gpu::GpuMat tested(grey);\r
+    int count = gpuClassifier.detectMultiScale(tested, buffer, gpu_rects);\r
+\r
+    cv::Mat gpu_f(gpu_rects);\r
+    int* gpu_faces = (int*)gpu_f.ptr();\r
+    for (int i = 0; i < count; i++)\r
+    {\r
+        cv::Rect r(gpu_faces[i * 4],gpu_faces[i * 4 + 1],gpu_faces[i * 4 + 2],gpu_faces[i * 4 + 3]);\r
+        cv::rectangle(markedImage, r , cv::Scalar(0, 0, 255, 255));\r
+    }\r
+}\r
+\r
+INSTANTIATE_TEST_CASE_P(GPU_ObjDetect, LBP_classify, testing::Combine(\r
+    ALL_DEVICES,\r
+    testing::Values<int>(0)\r
+    ));\r
+\r
 } // namespace\r