changed mean shift procedure a little

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 7d0074d..8bb7c2a 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/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, \r
             TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1));\r
 \r
+        //! Does mean shift procedure on GPU.\r
+        CV_EXPORTS void meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int sp, int sr, \r
+            TermCriteria criteria = TermCriteria(TermCriteria::MAX_ITER + TermCriteria::EPS, 5, 1));\r
+\r
         //! Does coloring of disparity image: [0..ndisp) -> [0..240, 1, 1] in HSV.\r
         //! Supported types of input disparity: CV_8U, CV_16S.\r
         //! Output disparity has CV_8UC4 type in BGRA format (alpha = 255).\r

diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu
index 1dd4616..938e2d1 100644 (file)
--- a/modules/gpu/src/cuda/imgproc.cu
+++ b/modules/gpu/src/cuda/imgproc.cu
@@ -163,70 +163,93 @@ namespace imgproc
 {\r
     texture<uchar4, 2> tex_meanshift;\r
 \r
-    extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, \r
-                                                 int sp, int sr, int maxIter, float eps )\r
+    __device__ short2 do_mean_shift(int x0, int y0, unsigned char* out, \r
+                                    int out_step, int cols, int rows, \r
+                                    int sp, int sr, int maxIter, float eps)\r
     {\r
-        int x0 = blockIdx.x * blockDim.x + threadIdx.x;\r
-        int y0 = blockIdx.y * blockDim.y + threadIdx.y;\r
+        int isr2 = sr*sr;\r
+        uchar4 c = tex2D(tex_meanshift, x0, y0 );\r
 \r
-        if( x0 < cols && y0 < rows )\r
+        // iterate meanshift procedure\r
+        for( int iter = 0; iter < maxIter; iter++ )\r
         {\r
-            int isr2 = sr*sr;\r
-            uchar4 c = tex2D(tex_meanshift, x0, y0 );\r
-            // iterate meanshift procedure\r
-            for( int iter = 0; iter < maxIter; iter++ )\r
+            int count = 0;\r
+            int s0 = 0, s1 = 0, s2 = 0, sx = 0, sy = 0;\r
+            float icount;\r
+\r
+            //mean shift: process pixels in window (p-sigmaSp)x(p+sigmaSp)\r
+            int minx = x0-sp;\r
+            int miny = y0-sp;\r
+            int maxx = x0+sp;\r
+            int maxy = y0+sp;\r
+\r
+            for( int y = miny; y <= maxy; y++)\r
             {\r
-                int count = 0;\r
-                int s0 = 0, s1 = 0, s2 = 0, sx = 0, sy = 0;\r
-                float icount;\r
-\r
-                //mean shift: process pixels in window (p-sigmaSp)x(p+sigmaSp)\r
-                int minx = x0-sp;\r
-                int miny = y0-sp;\r
-                int maxx = x0+sp;\r
-                int maxy = y0+sp;\r
-\r
-                for( int y = miny; y <= maxy; y++)\r
-                {\r
-                    int rowCount = 0;\r
-                    for( int x = minx; x <= maxx; x++ )\r
-                    {                    \r
-                        uchar4 t = tex2D( tex_meanshift, x, y );\r
-\r
-                        int norm2 = (t.x - c.x) * (t.x - c.x) + (t.y - c.y) * (t.y - c.y) + (t.z - c.z) * (t.z - c.z);\r
-                        if( norm2 <= isr2 )\r
-                        {\r
-                            s0 += t.x; s1 += t.y; s2 += t.z;\r
-                            sx += x; rowCount++;\r
-                        }\r
+                int rowCount = 0;\r
+                for( int x = minx; x <= maxx; x++ )\r
+                {                    \r
+                    uchar4 t = tex2D( tex_meanshift, x, y );\r
+\r
+                    int norm2 = (t.x - c.x) * (t.x - c.x) + (t.y - c.y) * (t.y - c.y) + (t.z - c.z) * (t.z - c.z);\r
+                    if( norm2 <= isr2 )\r
+                    {\r
+                        s0 += t.x; s1 += t.y; s2 += t.z;\r
+                        sx += x; rowCount++;\r
                     }\r
-                    count += rowCount;\r
-                    sy += y*rowCount;\r
                 }\r
+                count += rowCount;\r
+                sy += y*rowCount;\r
+            }\r
 \r
-                if( count == 0 )\r
-                    break;\r
+            if( count == 0 )\r
+                break;\r
 \r
-                icount = 1.f/count;\r
-                int x1 = __float2int_rz(sx*icount);\r
-                int y1 = __float2int_rz(sy*icount);\r
-                s0 = __float2int_rz(s0*icount);\r
-                s1 = __float2int_rz(s1*icount);\r
-                s2 = __float2int_rz(s2*icount);\r
+            icount = 1.f/count;\r
+            int x1 = __float2int_rz(sx*icount);\r
+            int y1 = __float2int_rz(sy*icount);\r
+            s0 = __float2int_rz(s0*icount);\r
+            s1 = __float2int_rz(s1*icount);\r
+            s2 = __float2int_rz(s2*icount);\r
 \r
-                int norm2 = (s0 - c.x) * (s0 - c.x) + (s1 - c.y) * (s1 - c.y) + (s2 - c.z) * (s2 - c.z);\r
+            int norm2 = (s0 - c.x) * (s0 - c.x) + (s1 - c.y) * (s1 - c.y) + (s2 - c.z) * (s2 - c.z);\r
 \r
-                bool stopFlag = (x0 == x1 && y0 == y1) || (abs(x1-x0) + abs(y1-y0) + norm2 <= eps);\r
+            bool stopFlag = (x0 == x1 && y0 == y1) || (abs(x1-x0) + abs(y1-y0) + norm2 <= eps);\r
 \r
-                x0 = x1; y0 = y1;\r
-                c.x = s0; c.y = s1; c.z = s2;\r
+            x0 = x1; y0 = y1;\r
+            c.x = s0; c.y = s1; c.z = s2;\r
 \r
-                if( stopFlag )\r
-                    break;\r
-            }\r
+            if( stopFlag )\r
+                break;\r
+        }\r
+\r
+        int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 4 * sizeof(uchar);\r
+        *(uchar4*)(out + base) = c;\r
+\r
+        return make_short2((short)x0, (short)y0);\r
+    }\r
+\r
+    extern "C" __global__ void meanshift_kernel( unsigned char* out, int out_step, int cols, int rows, \r
+                                                 int sp, int sr, int maxIter, float eps )\r
+    {\r
+        int x0 = blockIdx.x * blockDim.x + threadIdx.x;\r
+        int y0 = blockIdx.y * blockDim.y + threadIdx.y;\r
+\r
+        if( x0 < cols && y0 < rows )\r
+            do_mean_shift(x0, y0, out, out_step, cols, rows, sp, sr, maxIter, eps);\r
+    }\r
+\r
+    extern "C" __global__ void meanshiftproc_kernel( unsigned char* outr, int outrstep, \r
+                                                 unsigned char* outsp, int outspstep, \r
+                                                 int cols, int rows, \r
+                                                 int sp, int sr, int maxIter, float eps )\r
+    {\r
+        int x0 = blockIdx.x * blockDim.x + threadIdx.x;\r
+        int y0 = blockIdx.y * blockDim.y + threadIdx.y;\r
 \r
-            int base = (blockIdx.y * blockDim.y + threadIdx.y) * out_step + (blockIdx.x * blockDim.x + threadIdx.x) * 4 * sizeof(uchar);\r
-            *(uchar4*)(out + base) = c;\r
+        if( x0 < cols && y0 < rows )\r
+        {            \r
+            int basesp = (blockIdx.y * blockDim.y + threadIdx.y) * outspstep + (blockIdx.x * blockDim.x + threadIdx.x) * 2 * sizeof(short);\r
+            *(short2*)(outsp + basesp) = do_mean_shift(x0, y0, outr, outrstep, cols, rows, sp, sr, maxIter, eps);\r
         }\r
     }\r
 }\r
@@ -247,6 +270,20 @@ namespace cv { namespace gpu { namespace improc
         cudaSafeCall( cudaThreadSynchronize() );\r
         cudaSafeCall( cudaUnbindTexture( imgproc::tex_meanshift ) );        \r
     }\r
+    extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps) \r
+    {\r
+        dim3 grid(1, 1, 1);\r
+        dim3 threads(32, 16, 1);\r
+        grid.x = divUp(src.cols, threads.x);\r
+        grid.y = divUp(src.rows, threads.y);\r
+\r
+        cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar4>();\r
+        cudaSafeCall( cudaBindTexture2D( 0, imgproc::tex_meanshift, src.ptr, desc, src.cols, src.rows, src.step ) );\r
+\r
+        imgproc::meanshiftproc_kernel<<< grid, threads >>>( dstr.ptr, dstr.step, dstsp.ptr, dstsp.step, dstr.cols, dstr.rows, sp, sr, maxIter, eps );\r
+        cudaSafeCall( cudaThreadSynchronize() );\r
+        cudaSafeCall( cudaUnbindTexture( imgproc::tex_meanshift ) );        \r
+    }\r
 }}}\r
 \r
 /////////////////////////////////// drawColorDisp ///////////////////////////////////////////////\r

diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp
index 0badf8b..33bf5b5 100644 (file)
--- a/modules/gpu/src/imgproc_gpu.cpp
+++ b/modules/gpu/src/imgproc_gpu.cpp
@@ -49,6 +49,7 @@ using namespace cv::gpu;
 \r
 void cv::gpu::remap(const GpuMat&, GpuMat&, const GpuMat&, const GpuMat&){ throw_nogpu(); }\r
 void cv::gpu::meanShiftFiltering(const GpuMat&, GpuMat&, int, int, TermCriteria) { throw_nogpu(); }\r
+void cv::gpu::meanShiftProc(const GpuMat&, GpuMat&, GpuMat&, int, int, TermCriteria) { throw_nogpu(); }\r
 void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int) { throw_nogpu(); }\r
 void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, const Stream&) { throw_nogpu(); }\r
 void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&) { throw_nogpu(); }\r
@@ -74,6 +75,7 @@ namespace cv { namespace gpu
         void remap_gpu_3c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);\r
 \r
         extern "C" void meanShiftFiltering_gpu(const DevMem2D& src, DevMem2D dst, int sp, int sr, int maxIter, float eps);\r
+        extern "C" void meanShiftProc_gpu(const DevMem2D& src, DevMem2D dstr, DevMem2D dstsp, int sp, int sr, int maxIter, float eps);\r
 \r
         void drawColorDisp_gpu(const DevMem2D& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream);\r
         void drawColorDisp_gpu(const DevMem2D_<short>& src, const DevMem2D& dst, int ndisp, const cudaStream_t& stream);\r
@@ -164,6 +166,33 @@ void cv::gpu::meanShiftFiltering(const GpuMat& src, GpuMat& dst, int sp, int sr,
 }\r
 \r
 ////////////////////////////////////////////////////////////////////////\r
+// meanShiftProc_GPU\r
+\r
+void cv::gpu::meanShiftProc(const GpuMat& src, GpuMat& dstr, GpuMat& dstsp, int sp, int sr, TermCriteria criteria)\r
+{       \r
+    if( src.empty() )\r
+        CV_Error( CV_StsBadArg, "The input image is empty" );\r
+\r
+    if( src.depth() != CV_8U || src.channels() != 4 )\r
+        CV_Error( CV_StsUnsupportedFormat, "Only 8-bit, 4-channel images are supported" );\r
+\r
+    dstr.create( src.size(), CV_8UC4 );\r
+    dstsp.create( src.size(), CV_16SC2 );\r
+    \r
+    if( !(criteria.type & TermCriteria::MAX_ITER) )\r
+        criteria.maxCount = 5;\r
+    \r
+    int maxIter = std::min(std::max(criteria.maxCount, 1), 100);\r
+    \r
+    float eps;\r
+    if( !(criteria.type & TermCriteria::EPS) )\r
+        eps = 1.f;\r
+    eps = (float)std::max(criteria.epsilon, 0.0);        \r
+\r
+    improc::meanShiftProc_gpu(src, dstr, dstsp, sp, sr, maxIter, eps);    \r
+}\r
+\r
+////////////////////////////////////////////////////////////////////////\r
 // drawColorDisp\r
 \r
 namespace\r