minor refactoring of GPU module and GPU tests

added gpu compare version for CMP_NE operation

modules/gpu/include/opencv2/gpu/gpu.hpp

@@ -364,7 +364,7 @@ namespace cv
 
         //! applies fixed threshold to the image. 
         //! Now supports only THRESH_TRUNC threshold type and one channels float source.
-        CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh, double maxVal, int thresholdType);
+        CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh);
 
         //! compares elements of two arrays (c = a <cmpop> b)
         //! Now doesn't support CMP_NE.

modules/gpu/src/cuda/matrix_operations.cu

@@ -255,6 +255,24 @@ namespace mat_operators
         }
     }
 
+    ///////////////////////////////////////////////////////////////////////////
+    /////////////////////////////// compare_ne ////////////////////////////////
+    ///////////////////////////////////////////////////////////////////////////
+
+    template <typename T>
+    __global__ void kernel_compare_ne(uchar* src1, size_t src1_step, uchar* src2, size_t src2_step, uchar* dst, size_t dst_step, int cols, int rows)
+    {
+        const size_t x = threadIdx.x + blockIdx.x * blockDim.x;
+        const size_t y = threadIdx.y + blockIdx.y * blockDim.y;
+
+        if (x < cols && y < rows)
+        {
+            T src1_pix = ((T*)(src1 + y * src1_step))[x];
+            T src2_pix = ((T*)(src2 + y * src2_step))[x];
+            uchar res = (uchar)(src1_pix != src2_pix) * 255;
+            ((dst + y * dst_step))[x] = res;
+        }
+    }
 } // namespace mat_operators
 
 namespace cv
@@ -460,6 +478,28 @@ namespace cv
                     cv::gpu::error("Unsupported convert operation", __FILE__, __LINE__);
                 func(src, dst, src.cols * channels, src.rows, alpha, beta, stream);
             }
-        } // namespace impl
+
+            ///////////////////////////////////////////////////////////////////////////
+            /////////////////////////////// compare_ne ////////////////////////////////
+            ///////////////////////////////////////////////////////////////////////////
+
+            void compare_ne_8u(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst)
+            {
+                dim3 block(32, 8);
+                dim3 grid(divUp(src1.cols, block.x), divUp(src1.rows, block.y));
+
+                mat_operators::kernel_compare_ne<uint><<<grid, block>>>(src1.ptr, src1.step, src2.ptr, src2.step, dst.ptr, dst.step, src1.cols, src1.rows);
+                cudaSafeCall( cudaThreadSynchronize() );
+            }
+
+            void compare_ne_32f(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst)
+            {
+                dim3 block(32, 8);
+                dim3 grid(divUp(src1.cols, block.x), divUp(src1.rows, block.y));
+
+                mat_operators::kernel_compare_ne<float><<<grid, block>>>(src1.ptr, src1.step, src2.ptr, src2.step, dst.ptr, dst.step, src1.cols, src1.rows);
+                cudaSafeCall( cudaThreadSynchronize() );
+            }
+        } // namespace matrix_operations
     } // namespace gpu
 } // namespace cv

modules/gpu/src/filtering_npp.cpp

@@ -48,9 +48,9 @@ using namespace cv::gpu;
 
 #if !defined (HAVE_CUDA)
 
-void cv::gpu::erode( const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor, int iterations) { throw_nogpu(); }
-void cv::gpu::dilate( const GpuMat& src, GpuMat& dst, const Mat& kernel, Point anchor, int iterations) { throw_nogpu(); }
-void morphologyEx( const GpuMat& src, GpuMat& dst, int op, const Mat& kernel, Point anchor, int iterations) { throw_nogpu(); }
+void cv::gpu::erode( const GpuMat&, GpuMat&, const Mat&, Point, int) { throw_nogpu(); }
+void cv::gpu::dilate( const GpuMat&, GpuMat&, const Mat&, Point, int) { throw_nogpu(); }
+void morphologyEx( const GpuMat&, GpuMat&, int, const Mat&, Point, int) { throw_nogpu(); }
 
 #else
 

tests/gpu/src/meanshift.cpp

@@ -60,21 +60,23 @@ CV_GpuMeanShiftTest::CV_GpuMeanShiftTest(): CvTest( "GPU-MeanShift", "MeanShift"
 
 void CV_GpuMeanShiftTest::run(int)
 {
-        int spatialRad = 30;
-        int colorRad = 30;
+    int spatialRad = 30;
+    int colorRad = 30;
 
-        cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "meanshift/cones.png");
-        cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result.png");
+    cv::Mat img = cv::imread(std::string(ts->get_data_path()) + "meanshift/cones.png");
+    cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "meanshift/con_result.png");
 
-        if (img.empty() || img_template.empty())
-        {
-            ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
-            return;
-        }
+    if (img.empty() || img_template.empty())
+    {
+        ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
+        return;
+    }
 
-        cv::Mat rgba;
-        cvtColor(img, rgba, CV_BGR2BGRA);
+    cv::Mat rgba;
+    cvtColor(img, rgba, CV_BGR2BGRA);
 
+    try
+    {
         cv::gpu::GpuMat res;
         cv::gpu::meanShiftFiltering( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad );
         if (res.type() != CV_8UC4)
@@ -98,15 +100,27 @@ void CV_GpuMeanShiftTest::run(int)
                 {
                     const uchar& ch1 = res_line[result.channels()*i + k];
                     const uchar& ch2 = ref_line[img_template.channels()*i + k];
-                    uchar diff = abs(ch1 - ch2);
+                    uchar diff = static_cast<uchar>(abs(ch1 - ch2));
                     if (maxDiff < diff)
                         maxDiff = diff;
                 }
             }
         }
         if (maxDiff > 0) 
+        {
             ts->printf(CvTS::CONSOLE, "\nMeanShift maxDiff = %d\n", maxDiff);
-        ts->set_failed_test_info((maxDiff == 0) ? CvTS::OK : CvTS::FAIL_GENERIC);
+            ts->set_failed_test_info(CvTS::FAIL_GENERIC);
+            return;
+        }
+    }
+    catch(const cv::Exception& e)
+    {
+        if (!check_and_treat_gpu_exception(e, ts))
+            throw;
+        return;
+    }
+
+    ts->set_failed_test_info(CvTS::OK);
 }
 
-CV_GpuMeanShiftTest CV_GpuMeanShift_test;
\ No newline at end of file
+CV_GpuMeanShiftTest CV_GpuMeanShift_test;

tests/gpu/src/morf_filters.cpp

@@ -81,7 +81,7 @@ protected:
         if (res < std::numeric_limits<double>::epsilon())
             return CvTS::OK;
         
-        ts->printf(CvTS::LOG, "\nNorm: %f\n", res);
+        ts->printf(CvTS::CONSOLE, "\nNorm: %f\n", res);
         return CvTS::FAIL_GENERIC;
     }
 };
@@ -116,7 +116,8 @@ void CV_GpuNppMorphogyTest::run( int )
     catch(const cv::Exception& e)
     {
         if (!check_and_treat_gpu_exception(e, ts))
-            throw;        
+            throw;
+        return;
     }
 
     ts->set_failed_test_info(CvTS::OK);
@@ -174,7 +175,6 @@ protected:
 
 CV_GpuDilateTest CV_GpuDilate_test;
 
-
 ////////////////////////////////////////////////////////////////////////////////
 // Dilate
 class CV_GpuMorphExTest : public CV_GpuNppMorphogyTest

tests/gpu/src/operator_async_call.cpp

@@ -143,7 +143,16 @@ void CV_GpuMatAsyncCallTest::run( int /* start_from */)
     Mat cpumat(rows, cols, CV_8U);
     cpumat.setTo(Scalar::all(127));
 
-    is_test_good &= compare_matrix(cpumat);
+    try
+    {
+        is_test_good &= compare_matrix(cpumat);
+    }
+    catch(cv::Exception& e)
+    {
+        if (!check_and_treat_gpu_exception(e, ts))
+            throw; 
+        return;
+    }
 
     if (is_test_good == true)
         ts->set_failed_test_info(CvTS::OK);
@@ -151,4 +160,4 @@ void CV_GpuMatAsyncCallTest::run( int /* start_from */)
         ts->set_failed_test_info(CvTS::FAIL_GENERIC);
 }
 
-CV_GpuMatAsyncCallTest CV_GpuMatAsyncCall_test;
+//CV_GpuMatAsyncCallTest CV_GpuMatAsyncCall_test;

tests/gpu/src/operator_convert_to.cpp

@@ -115,10 +115,12 @@ void CV_GpuMatOpConvertToTest::run(int /* start_from */)
     }
     catch(cv::Exception& e)
     {
-        ts->printf(CvTS::CONSOLE, "\nERROR: %s\n", e.what());
+        if (!check_and_treat_gpu_exception(e, ts))
+            throw; 
+        return;
     }
+
     ts->set_failed_test_info(passed ? CvTS::OK : CvTS::FAIL_GENERIC);
 }
 
 CV_GpuMatOpConvertToTest CV_GpuMatOpConvertToTest_test;
-

tests/gpu/src/operator_copy_to.cpp

@@ -136,14 +136,23 @@ void CV_GpuMatOpCopyToTest::run( int /* start_from */)
 {
     bool is_test_good = true;
 
-    for (int i = 0 ; i < 7; i++)
+    try
     {
-        Mat cpumat(rows, cols, i);
-        cpumat.setTo(Scalar::all(127));
+        for (int i = 0 ; i < 7; i++)
+        {
+            Mat cpumat(rows, cols, i);
+            cpumat.setTo(Scalar::all(127));
 
-        GpuMat gpumat(cpumat);
+            GpuMat gpumat(cpumat);
 
-        is_test_good &= compare_matrix(cpumat, gpumat);
+            is_test_good &= compare_matrix(cpumat, gpumat);
+        }
+    }
+    catch(const cv::Exception& e)
+    {
+        if (!check_and_treat_gpu_exception(e, ts))
+            throw;
+        return;
     }
 
     if (is_test_good == true)

tests/gpu/src/operator_set_to.cpp

@@ -132,11 +132,20 @@ void CV_GpuMatOpSetToTest::run( int /* start_from */)
 {
     bool is_test_good = true;
 
-    for (int i = 0; i < 7; i++)
+    try
     {
-        Mat cpumat(rows, cols, i, Scalar::all(0));
-        GpuMat gpumat(cpumat);
-        is_test_good &= compare_matrix(cpumat, gpumat);
+        for (int i = 0; i < 7; i++)
+        {
+            Mat cpumat(rows, cols, i, Scalar::all(0));
+            GpuMat gpumat(cpumat);
+            is_test_good &= compare_matrix(cpumat, gpumat);
+        }
+    }
+    catch(const cv::Exception& e)
+    {
+        if (!check_and_treat_gpu_exception(e, ts))
+            throw;
+        return;
     }
 
     if (is_test_good == true)

tests/gpu/src/stereo_bm.cpp

@@ -70,18 +70,30 @@ void CV_GpuStereoBMTest::run(int )
         return;
     }
 
-	cv::gpu::GpuMat disp;
-	cv::gpu::StereoBM_GPU bm(0, 128, 19);
-	bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
+    try
+    {
+	    cv::gpu::GpuMat disp;
+	    cv::gpu::StereoBM_GPU bm(0, 128, 19);
+	    bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
 
-	disp.convertTo(disp, img_reference.type());
-	double norm = cv::norm(disp, img_reference, cv::NORM_INF);
+	    disp.convertTo(disp, img_reference.type());
+	    double norm = cv::norm(disp, img_reference, cv::NORM_INF);
 
-	if (norm >= 100) 
-        ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm);
-	ts->set_failed_test_info((norm < 100) ? CvTS::OK : CvTS::FAIL_GENERIC);
-}
+	    if (norm >= 100) 
+        {
+            ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm);
+	        ts->set_failed_test_info(CvTS::FAIL_GENERIC);
+            return;
+        }
+    }
+    catch(const cv::Exception& e)
+    {
+        if (!check_and_treat_gpu_exception(e, ts))
+            throw;
+        return;
+    }
 
+    ts->set_failed_test_info(CvTS::OK);
+}
 
 CV_GpuStereoBMTest CV_GpuStereoBM_test;
-

tests/gpu/src/stereo_bm_async.cpp

@@ -74,24 +74,37 @@ void CV_GpuMatAsyncCallStereoBMTest::run( int /* start_from */)
         return;
     }
 
-	cv::gpu::GpuMat disp;
-	cv::gpu::StereoBM_GPU bm(0, 128, 19);
-
-	cv::gpu::Stream stream;
-
-	for (size_t i = 0; i < 50; i++)
-	{
-		bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp, stream);
-	}
-
-	stream.waitForCompletion();
-	disp.convertTo(disp, img_reference.type());
-	double norm = cv::norm(disp, img_reference, cv::NORM_INF);
-
-	if (norm >= 100) 
-        ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm);
-	ts->set_failed_test_info((norm < 100) ? CvTS::OK : CvTS::FAIL_GENERIC);
+    try
+    {
+	    cv::gpu::GpuMat disp;
+	    cv::gpu::StereoBM_GPU bm(0, 128, 19);
+
+	    cv::gpu::Stream stream;
+
+	    for (size_t i = 0; i < 50; i++)
+	    {
+		    bm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp, stream);
+	    }
+
+	    stream.waitForCompletion();
+	    disp.convertTo(disp, img_reference.type());
+	    double norm = cv::norm(disp, img_reference, cv::NORM_INF);
+
+	    if (norm >= 100) 
+        {
+            ts->printf(CvTS::CONSOLE, "\nStereoBM norm = %f\n", norm);
+	        ts->set_failed_test_info(CvTS::FAIL_GENERIC);
+            return;
+        }
+    }
+    catch(const cv::Exception& e)
+    {
+        if (!check_and_treat_gpu_exception(e, ts))
+            throw;
+        return;
+    }
 
+    ts->set_failed_test_info(CvTS::OK);
 }
 
 CV_GpuMatAsyncCallStereoBMTest CV_GpuMatAsyncCallStereoBMTest_test;

tests/gpu/src/stereo_bp.cpp

@@ -69,20 +69,33 @@ void CV_GpuStereoBPTest::run(int )
         return;
     }
 
-    cv::gpu::GpuMat disp;
-    cv::gpu::StereoBeliefPropagation bpm(64, 8, 2, 25, 0.1f, 15, 1, CV_16S);
+    try
+    {
+        cv::gpu::GpuMat disp;
+        cv::gpu::StereoBeliefPropagation bpm(64, 8, 2, 25, 0.1f, 15, 1, CV_16S);
 
-    bpm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
+        bpm(cv::gpu::GpuMat(img_l), cv::gpu::GpuMat(img_r), disp);
 
-    //cv::imwrite(std::string(ts->get_data_path()) + "stereobp/aloe-disp.png", disp);
+        //cv::imwrite(std::string(ts->get_data_path()) + "stereobp/aloe-disp.png", disp);
 
-    disp.convertTo(disp, img_template.type());
+        disp.convertTo(disp, img_template.type());
 
-    double norm = cv::norm(disp, img_template, cv::NORM_INF);
-	if (norm >= 0.5) 
-        ts->printf(CvTS::CONSOLE, "\nStereoBP norm = %f\n", norm);
-    ts->set_failed_test_info((norm < 0.5) ? CvTS::OK : CvTS::FAIL_GENERIC);
-}
+        double norm = cv::norm(disp, img_template, cv::NORM_INF);
+	    if (norm >= 0.5) 
+        {
+            ts->printf(CvTS::CONSOLE, "\nStereoBP norm = %f\n", norm);
+            ts->set_failed_test_info(CvTS::FAIL_GENERIC);
+            return;
+        }
+    }
+    catch(const cv::Exception& e)
+    {
+        if (!check_and_treat_gpu_exception(e, ts))
+            throw;
+        return;
+    }
 
+    ts->set_failed_test_info(CvTS::OK);
+}
 
-CV_GpuStereoBPTest CV_GpuStereoBP_test;
\ No newline at end of file
+CV_GpuStereoBPTest CV_GpuStereoBP_test;

tests/gpu/src/stereo_csbp.cpp

@@ -59,16 +59,18 @@ CV_GpuStereoCSBPTest::CV_GpuStereoCSBPTest(): CvTest( "GPU-StereoCSBP", "Constan
 
 void CV_GpuStereoCSBPTest::run(int )
 {
-        cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-L.png");
-        cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-R.png");
-        cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-disp.png", 0);
+    cv::Mat img_l = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-L.png");
+    cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-R.png");
+    cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "csstereobp/aloe-disp.png", 0);
 
-        if (img_l.empty() || img_r.empty() || img_template.empty())
-        {
-            ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
-            return;
-        }
+    if (img_l.empty() || img_r.empty() || img_template.empty())
+    {
+        ts->set_failed_test_info(CvTS::FAIL_MISSING_TEST_DATA);
+        return;
+    }
 
+    try
+    {
         cv::gpu::GpuMat disp;
         cv::gpu::StereoConstantSpaceBP bpm(128, 16, 4, 4);
 
@@ -79,9 +81,21 @@ void CV_GpuStereoCSBPTest::run(int )
         disp.convertTo(disp, img_template.type());
 
         double norm = cv::norm(disp, img_template, cv::NORM_INF);
-		if (norm >= 0.5) 
+	    if (norm >= 0.5) 
+        {
             ts->printf(CvTS::CONSOLE, "\nConstantSpaceStereoBP norm = %f\n", norm);
-        ts->set_failed_test_info((norm < 0.5) ? CvTS::OK : CvTS::FAIL_GENERIC);
+            ts->set_failed_test_info(CvTS::FAIL_GENERIC);
+            return;
+        }
+    }
+    catch(const cv::Exception& e)
+    {
+        if (!check_and_treat_gpu_exception(e, ts))
+            throw;
+        return;
+    }
+
+    ts->set_failed_test_info(CvTS::OK);
 }
 
-CV_GpuStereoCSBPTest CV_GpuCSStereoBP_test;
\ No newline at end of file
+CV_GpuStereoCSBPTest CV_GpuCSStereoBP_test;