meanShiftFiltering_GPU output parameters changed to CV_8UC4. This is a start for moving from 3 channel to C4 images within GPU module.

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

@@ -347,7 +347,8 @@ namespace cv
         CV_EXPORTS void remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, GpuMat& dst);
 
         // Does mean shift filtering on GPU.
-        CV_EXPORTS void meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int sr, TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1));
+        CV_EXPORTS void meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int sr, 
+            TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1));
 
         // Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV.
         // Supported types of input disparity: CV_8U, CV_16S.

modules/gpu/src/cuda/imgproc.cu

@@ -163,7 +163,8 @@ namespace imgproc
 {
     texture<uchar4, 2> tex_meanshift;
 
-    extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, int sp, int sr, int maxIter, float eps )
+    extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, 
+                                                 int sp, int sr, int maxIter, float eps )
     {
         int x0 = blockIdx.x * blockDim.x + threadIdx.x;
         int y0 = blockIdx.y * blockDim.y + threadIdx.y;
@@ -224,10 +225,8 @@ namespace imgproc
                     break;
             }
 
-            int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 3 * sizeof(uchar);
+            int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 4 * sizeof(uchar);
-            out[base+0] = c.x;
+            *(uchar4*)(out + base) = c;
-            out[base+1] = c.y;
-            out[base+2] = c.z;
         }
     }
 }

modules/gpu/src/imgproc_gpu.cpp

@@ -119,18 +119,18 @@ void cv::gpu::meanShiftFiltering_GPU(const GpuMat& src, GpuMat& dst, int sp, int
     if( src.depth() != CV_8U || src.channels() != 4 )
         CV_Error( CV_StsUnsupportedFormat, "Only 8-bit, 4-channel images are supported" );
 
-    dst.create( src.size(), CV_8UC3 );
+    dst.create( src.size(), CV_8UC4 );
     
-    float eps;
     if( !(criteria.type & TermCriteria::MAX_ITER) )
         criteria.maxCount = 5;
     
     int maxIter = std::min(std::max(criteria.maxCount, 1), 100);
     
+    float eps;
     if( !(criteria.type & TermCriteria::EPS) )
         eps = 1.f;
-
     eps = (float)std::max(criteria.epsilon, 0.0);        
+
     impl::meanShiftFiltering_gpu(src, dst, sp, sr, maxIter, eps);    
 }
 

tests/gpu/src/meanshift.cpp

@@ -56,7 +56,7 @@ class CV_GpuMeanShift : public CvTest
 
 CV_GpuMeanShift::CV_GpuMeanShift(): CvTest( "GPU-MeanShift", "MeanShift" ){}
 
-void CV_GpuMeanShift::run(int )
+void CV_GpuMeanShift::run(int)
 {
         int spatialRad = 30;
         int colorRad = 30;
@@ -64,19 +64,47 @@ void CV_GpuMeanShift::run(int )
         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;
+        }
+
         cv::Mat rgba;
         cvtColor(img, rgba, CV_BGR2BGRA);
 
         cv::gpu::GpuMat res;
-
         cv::gpu::meanShiftFiltering_GPU( cv::gpu::GpuMat(rgba), res, spatialRad, colorRad );
+        if (res.type() != CV_8UC4)
+        {
+            ts->set_failed_test_info(CvTS::FAIL_INVALID_OUTPUT);
+            return;
+        }
 
-        res.convertTo(res, img_template.type());
+        cv::Mat result;
+        res.download(result);
 
-        double norm = cv::norm(res, img_template, cv::NORM_INF);
+        uchar maxDiff = 0;
-		if (norm >= 0.5) std::cout << "MeanShift norm = " << norm << std::endl;
+        for (int j = 0; j < result.rows; ++j)
-        ts->set_failed_test_info((norm < 0.5) ? CvTS::OK : CvTS::FAIL_GENERIC);
+        {
+            const uchar* res_line = result.ptr<uchar>(j);
+            const uchar* ref_line = img_template.ptr<uchar>(j);
+
+            for (int i = 0; i < result.cols; ++i)
+            {
+                for (int k = 0; k < 3; ++k)
+                {
+                    const uchar& ch1 = res_line[result.channels()*i + k];
+                    const uchar& ch2 = ref_line[img_template.channels()*i + k];
+                    uchar diff = 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);
 }
 
-
 CV_GpuMeanShift CV_GpuMeanShift_test;

tests/gpu/src/stereo_bp.cpp

@@ -62,6 +62,12 @@ void CV_GpuStereoBP::run(int )
     cv::Mat img_r = cv::imread(std::string(ts->get_data_path()) + "stereobp/aloe-R.png");
     cv::Mat img_template = cv::imread(std::string(ts->get_data_path()) + "stereobp/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;
+    }
+
     cv::gpu::GpuMat disp;
     cv::gpu::StereoBeliefPropagation bpm(64, 8, 2, 25, 0.1f, 15, 1, CV_16S);