changed mean shift procedure a little

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

@@ -469,6 +469,10 @@ namespace cv
         CV_EXPORTS void meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr, 
             TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1));
 
+        //! Does mean shift procedure on GPU.
+        CV_EXPORTS void meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, 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.
         //! Output disparity has CV_8UC4 type in BGRA format (alpha = 255).

modules/gpu/src/cuda/imgproc.cu

@@ -163,16 +163,13 @@ namespace imgproc
 {
     texture<uchar4, 2> tex_meanshift;
 
-    extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, 
+    __device__ short2 do_mean_shift(int x0, int y0, 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;
-
-        if( x0 < cols && y0 < rows )
     {
         int isr2 = sr*sr;
         uchar4 c = tex2D(tex_meanshift, x0, y0 );
+
         // iterate meanshift procedure
         for( int iter = 0; iter < maxIter; iter++ )
         {
@@ -227,6 +224,32 @@ namespace imgproc
 
         int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 4 * sizeof(uchar);
         *(uchar4*)(out + base) = c;
+
+        return make_short2((short)x0, (short)y0);
+    }
+
+    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;
+
+        if( x0 < cols && y0 < rows )
+            do_mean_shift(x0, y0, out, out_step, cols, rows, sp, sr, maxIter, eps);
+    }
+
+    extern "C" __global__ void meanshiftproc_kernel( unsigned char* outr, int outrstep, 
+                                                 unsigned char* outsp, int outspstep, 
+                                                 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;
+
+        if( x0 < cols && y0 < rows )
+        {            
+            int basesp = (blockIdx.y * blockDim.y + threadIdx.y) * outspstep + (blockIdx.x * blockDim.x + threadIdx.x) * 2 * sizeof(short);
+            *(short2*)(outsp + basesp) = do_mean_shift(x0, y0, outr, outrstep, cols, rows, sp, sr, maxIter, eps);
         }
     }
 }
@@ -247,6 +270,20 @@ namespace cv { namespace gpu { namespace improc
         cudaSafeCall( cudaThreadSynchronize() );
         cudaSafeCall( cudaUnbindTexture( imgproc::tex_meanshift ) );        
     }
+    extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps) 
+    {
+        dim3 grid(1, 1, 1);
+        dim3 threads(32, 16, 1);
+        grid.x = divUp(src.cols, threads.x);
+        grid.y = divUp(src.rows, threads.y);
+
+        cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar4>();
+        cudaSafeCall( cudaBindTexture2D( 0, imgproc::tex_meanshift, src.ptr, desc, src.cols, src.rows, src.step ) );
+
+        imgproc::meanshiftproc_kernel<<< grid, threads >>>( dstr.ptr, dstr.step, dstsp.ptr, dstsp.step, dstr.cols, dstr.rows, sp, sr, maxIter, eps );
+        cudaSafeCall( cudaThreadSynchronize() );
+        cudaSafeCall( cudaUnbindTexture( imgproc::tex_meanshift ) );        
+    }
 }}}
 
 /////////////////////////////////// drawColorDisp ///////////////////////////////////////////////

modules/gpu/src/imgproc_gpu.cpp

@@ -49,6 +49,7 @@ using namespace cv::gpu;
 
 void cv::gpu::remap(const GpuMat&, GpuMat&, const GpuMat&, const GpuMat&){ throw_nogpu(); }
 void cv::gpu::meanShiftFiltering(const GpuMat&, GpuMat&, int, int, TermCriteria) { throw_nogpu(); }
+void cv::gpu::meanShiftProc(const GpuMat&, GpuMat&, GpuMat&, int, int, TermCriteria) { throw_nogpu(); }
 void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int) { throw_nogpu(); }
 void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, const Stream&) { throw_nogpu(); }
 void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&) { throw_nogpu(); }
@@ -74,6 +75,7 @@ namespace cv { namespace gpu
         void remap_gpu_3c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
 
         extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps);
+        extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps);
 
         void drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream);
         void drawColorDisp_gpu(const DevMem2D_<short>& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream);
@@ -163,6 +165,33 @@ void cv::gpu::meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr,
     improc::meanShiftFiltering_gpu(src, dst, sp, sr, maxIter, eps);    
 }
 
+////////////////////////////////////////////////////////////////////////
+// meanShiftProc_GPU
+
+void cv::gpu::meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int sp, int sr, TermCriteria criteria)
+{       
+    if( src.empty() )
+        CV_Error( CV_StsBadArg, "The input image is empty" );
+
+    if( src.depth() != CV_8U || src.channels() != 4 )
+        CV_Error( CV_StsUnsupportedFormat, "Only 8-bit, 4-channel images are supported" );
+
+    dstr.create( src.size(), CV_8UC4 );
+    dstsp.create( src.size(), CV_16SC2 );
+    
+    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);        
+
+    improc::meanShiftProc_gpu(src, dstr, dstsp, sp, sr, maxIter, eps);    
+}
+
 ////////////////////////////////////////////////////////////////////////
 // drawColorDisp