added performance test for LBP classifier.

modules/gpu/perf/perf_objdetect.cpp

@@ -57,4 +57,31 @@ GPU_PERF_TEST_1(HaarClassifier, cv::gpu::DeviceInfo)
 
 INSTANTIATE_TEST_CASE_P(ObjDetect, HaarClassifier, ALL_DEVICES);
 
+//===================== LBP cascade ==========================//
+GPU_PERF_TEST_1(LBPClassifier, cv::gpu::DeviceInfo)
+{
+    cv::gpu::DeviceInfo devInfo = GetParam();
+    cv::gpu::setDevice(devInfo.deviceID());
+
+    cv::Mat img_host = readImage("gpu/haarcascade/group_1_640x480_VGA.pgm", cv::IMREAD_GRAYSCALE);
+    ASSERT_FALSE(img_host.empty());
+
+    cv::gpu::CascadeClassifier_GPU_LBP cascade;
+
+    ASSERT_TRUE(cascade.load(perf::TestBase::getDataPath("gpu/lbpcascade/lbpcascade_frontalface.xml")));
+
+    cv::gpu::GpuMat img(img_host);
+        cv::gpu::GpuMat gpu_rects, buffer;
+
+    // cascade.detectMultiScale(img, objects_buffer);
+    cascade.detectMultiScale(img, buffer, gpu_rects);
+
+    TEST_CYCLE()
+    {
+        cascade.detectMultiScale(img, buffer, gpu_rects);
+    }
+}
+
+INSTANTIATE_TEST_CASE_P(ObjDetect, LBPClassifier, ALL_DEVICES);
+
 #endif

modules/gpu/perf_cpu/perf_objdetect.cpp

@@ -49,4 +49,26 @@ GPU_PERF_TEST_1(HaarClassifier, cv::gpu::DeviceInfo)
 
 INSTANTIATE_TEST_CASE_P(ObjDetect, HaarClassifier, ALL_DEVICES);
 
+//===================== LBP cascade ==========================//
+GPU_PERF_TEST_1(LBPClassifier, cv::gpu::DeviceInfo)
+{
+    cv::Mat img = readImage("gpu/haarcascade/group_1_640x480_VGA.pgm", cv::IMREAD_GRAYSCALE);
+    ASSERT_FALSE(img.empty());
+
+    cv::CascadeClassifier cascade;
+
+    ASSERT_TRUE(cascade.load(perf::TestBase::getDataPath("gpu/lbpcascade/lbpcascade_frontalface.xml")));
+
+    std::vector<cv::Rect> rects;
+
+    cascade.detectMultiScale(img, rects);
+
+    TEST_CYCLE()
+    {
+        cascade.detectMultiScale(img, rects);
+    }
+}
+
+INSTANTIATE_TEST_CASE_P(ObjDetect, LBPClassifier, ALL_DEVICES);
+
 #endif

modules/gpu/src/cascadeclassifier.cpp

@@ -350,7 +350,9 @@ int cv::gpu::CascadeClassifier_GPU_LBP::detectMultiScale(const GpuMat& image, Gp
 
     if (groupThreshold <= 0  || objects.empty())
         return 0;
-    return cv::gpu::device::lbp::connectedConmonents(candidates, groupThreshold, grouping_eps, dclassified);
+    cv::gpu::device::lbp::connectedConmonents(candidates, groupThreshold, grouping_eps, dclassified);
+    cudaSafeCall( cudaDeviceSynchronize() );
+    return *classified;
 }
 
 // ============ old fashioned haar cascade ==============================================//

modules/gpu/test/test_warp_affine.cpp

@@ -185,8 +185,6 @@ PARAM_TEST_CASE(WarpAffine, cv::gpu::DeviceInfo, cv::Size, MatType, Inverse, Int
     int borderType;
     bool useRoi;
 
-    cv::Mat M;
-
     virtual void SetUp()
     {
         devInfo = GET_PARAM(0);

modules/gpu/test/test_warp_perspective.cpp

@@ -185,8 +185,6 @@ PARAM_TEST_CASE(WarpPerspective, cv::gpu::DeviceInfo, cv::Size, MatType, Inverse
     int borderType;
     bool useRoi;
 
-    cv::Mat M;
-
     virtual void SetUp()
     {
         devInfo = GET_PARAM(0);