added remap version for 3-channels input images

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

@@ -415,7 +415,7 @@ namespace cv
             enum { DEFAULT_ITERS  = 5  };
             enum { DEFAULT_LEVELS = 5  };
 
-            static void estimateRecopmmendedParams( int width, int height, int & ndisp, int & iters, int & levels);
+            static void estimateRecommendedParams(int width, int height, int& ndisp, int& iters, int& levels);
 
             //! the default constructor
             explicit StereoBeliefPropagation(int ndisp  = DEFAULT_NDISP,
@@ -473,7 +473,7 @@ namespace cv
             enum { DEFAULT_LEVELS   = 4   };
             enum { DEFAULT_NR_PLANE = 4   };
 
-            static void estimateRecopmmendedParams( int width, int height, int & ndisp, int & iters, int & levels, int & nr_plane);
+            static void estimateRecommendedParams(int width, int height, int& ndisp, int& iters, int& levels, int& nr_plane);
 
             //! the default constructor
             explicit StereoConstantSpaceBP(int ndisp    = DEFAULT_NDISP,

modules/gpu/src/beliefpropagation_gpu.cpp

@@ -48,6 +48,8 @@ using namespace std;
 
 #if !defined (HAVE_CUDA)
 
+void cv::gpu::StereoBeliefPropagation::estimateRecommendedParams(int, int, int&, int&, int&) { throw_nogpu(); }
+
 cv::gpu::StereoBeliefPropagation::StereoBeliefPropagation(int, int, int, int) { throw_nogpu(); }
 cv::gpu::StereoBeliefPropagation::StereoBeliefPropagation(int, int, int, float, float, float, float, int) { throw_nogpu(); }
 
@@ -78,13 +80,13 @@ namespace
 }
 
 
-void cv::gpu::StereoBeliefPropagation::estimateRecopmmendedParams( int width, int height, int & ndisp, int & iters, int & levels)
+void cv::gpu::StereoBeliefPropagation::estimateRecommendedParams(int width, int height, int& ndisp, int& iters, int& levels)
 {
     ndisp = width / 4;
     if ((ndisp & 1) != 0) 
         ndisp++;
 
-    int mm =::max(width, height);
+    int mm = ::max(width, height);
     iters = mm / 100 + 2;
 
     levels = (int)(log(static_cast<double>(mm)) + 1) * 4 / 5;

modules/gpu/src/constantspacebp_gpu.cpp

@@ -48,6 +48,8 @@ using namespace std;
 
 #if !defined (HAVE_CUDA)
 
+void cv::gpu::StereoConstantSpaceBP::estimateRecommendedParams(int, int, int&, int&, int&, int&) { throw_nogpu(); }
+
 cv::gpu::StereoConstantSpaceBP::StereoConstantSpaceBP(int, int, int, int, int) { throw_nogpu(); }
 cv::gpu::StereoConstantSpaceBP::StereoConstantSpaceBP(int, int, int, int, float, float, float, float, int, int) { throw_nogpu(); }
 
@@ -105,10 +107,9 @@ namespace
     const float DEFAULT_DISC_SINGLE_JUMP = 10.0f;
 }
 
-void cv::gpu::StereoConstantSpaceBP::estimateRecopmmendedParams( int width, int height, int & ndisp, int & iters, int & levels, int &nr_plane)
+void cv::gpu::StereoConstantSpaceBP::estimateRecommendedParams(int width, int height, int& ndisp, int& iters, int& levels, int& nr_plane)
 {
     ndisp = (int) ((float) width / 3.14f);
-    if (ndisp & 1 != 0) ndisp++;
     if ((ndisp & 1) != 0) 
         ndisp++;
 

modules/gpu/src/cuda/imgproc.cu

@@ -44,13 +44,12 @@
 
 using namespace cv::gpu;
 
-
 /////////////////////////////////// Remap ///////////////////////////////////////////////
 namespace imgproc
 {
     texture<unsigned char, 2, cudaReadModeNormalizedFloat> tex_remap;
 
-    __global__ void kernel_remap(const float *mapx, const float *mapy, size_t map_step, unsigned char* out, size_t out_step, int width, int height)
+    __global__ void remap_1c(const float* mapx, const float* mapy, size_t map_step, uchar* out, size_t out_step, int width, int height)
     {    
         int x = blockDim.x * blockIdx.x + threadIdx.x;
         int y = blockDim.y * blockIdx.y + threadIdx.y;
@@ -65,27 +64,88 @@ namespace imgproc
         }
     }
 
+    __global__ void remap_3c(const uchar* src, size_t src_step, const float* mapx, const float* mapy, size_t map_step, 
+                             uchar* dst, size_t dst_step, int width, int height)
+    {    
+        const int x = blockDim.x * blockIdx.x + threadIdx.x;
+        const int y = blockDim.y * blockIdx.y + threadIdx.y;
+
+        if (x < width && y < height)
+        {
+            const int idx = y * (map_step >> 2) + x; /* map_step >> 2  <=> map_step / sizeof(float)*/
+
+            const float xcoo = mapx[idx];
+            const float ycoo = mapy[idx];
+            
+            uchar3 out = make_uchar3(0, 0, 0);
+
+            if (xcoo >= 0 && xcoo < width - 1 && ycoo >= 0 && ycoo < height - 1)
+            {
+                const int x1 = __float2int_rd(xcoo);
+                const int y1 = __float2int_rd(ycoo);
+                const int x2 = x1 + 1;
+                const int y2 = y1 + 1;
+                
+                uchar3 src_reg = *(uchar3*)(src + y1 * src_step + 3 * x1);
+                out.x += src_reg.x * (x2 - xcoo) * (y2 - ycoo);
+                out.y += src_reg.y * (x2 - xcoo) * (y2 - ycoo);
+                out.z += src_reg.z * (x2 - xcoo) * (y2 - ycoo);
+
+                src_reg = *(uchar3*)(src + y1 * src_step + 3 * x2);
+                
+                out.x += src_reg.x * (xcoo - x1) * (y2 - ycoo);
+                out.y += src_reg.y * (xcoo - x1) * (y2 - ycoo);
+                out.z += src_reg.z * (xcoo - x1) * (y2 - ycoo);
+
+                src_reg = *(uchar3*)(src + y2 * src_step + 3 * x1);
+                
+                out.x += src_reg.x * (x2 - xcoo) * (ycoo - y1);
+                out.y += src_reg.y * (x2 - xcoo) * (ycoo - y1);
+                out.z += src_reg.z * (x2 - xcoo) * (ycoo - y1);
+
+                src_reg = *(uchar3*)(src + y2 * src_step + 3 * x2);
+                
+                out.x += src_reg.x * (xcoo - x1) * (ycoo - y1);
+                out.y += src_reg.y * (xcoo - x1) * (ycoo - y1);
+                out.z += src_reg.z * (xcoo - x1) * (ycoo - y1);
+            }
+
+            *(uchar3*)(dst + y * dst_step + 3 * x) = out;
+        }
+    }
 }
 
 namespace cv { namespace gpu { namespace impl 
 {
-    extern "C" void remap_gpu(const DevMem2D& src, const DevMem2D_<float>& xmap, const DevMem2D_<float>& ymap, DevMem2D dst)
+    void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst)
     {
-        dim3 block(16, 16, 1);
+        dim3 threads(16, 16, 1);
         dim3 grid(1, 1, 1);
-        grid.x = divUp(dst.cols, block.x);
+        grid.x = divUp(dst.cols, threads.x);
-        grid.y = divUp(dst.rows, block.y);
+        grid.y = divUp(dst.rows, threads.y);
 
         imgproc::tex_remap.filterMode = cudaFilterModeLinear;	    
         imgproc::tex_remap.addressMode[0] = imgproc::tex_remap.addressMode[1] = cudaAddressModeWrap;
         cudaChannelFormatDesc desc = cudaCreateChannelDesc<unsigned char>();
-        cudaSafeCall( cudaBindTexture2D(0, imgproc::tex_remap, src.ptr, desc, dst.cols, dst.rows, src.step) );
+        cudaSafeCall( cudaBindTexture2D(0, imgproc::tex_remap, src.ptr, desc, src.cols, src.rows, src.step) );
 
-        imgproc::kernel_remap<<<grid, block>>>(xmap.ptr, ymap.ptr, xmap.step, dst.ptr, dst.step, dst.cols, dst.rows);
+        imgproc::remap_1c<<<grid, threads>>>(xmap.ptr, ymap.ptr, xmap.step, dst.ptr, dst.step, dst.cols, dst.rows);
 
         cudaSafeCall( cudaThreadSynchronize() );  
         cudaSafeCall( cudaUnbindTexture(imgproc::tex_remap) );
     }
+    
+    void remap_gpu_3c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst)
+    {
+        dim3 threads(32, 8, 1);
+        dim3 grid(1, 1, 1);
+        grid.x = divUp(dst.cols, threads.x);
+        grid.y = divUp(dst.rows, threads.y);
+
+        imgproc::remap_3c<<<grid, threads>>>(src.ptr, src.step, xmap.ptr, ymap.ptr, xmap.step, dst.ptr, dst.step, dst.cols, dst.rows);
+
+        cudaSafeCall( cudaThreadSynchronize() ); 
+    }
 }}}
 
 

modules/gpu/src/imgproc_gpu.cpp

@@ -60,7 +60,8 @@ namespace cv { namespace gpu
 { 
     namespace impl 
     {
-        extern "C" void remap_gpu(const DevMem2D& src, const DevMem2D_<float>& xmap, const DevMem2D_<float>& ymap, DevMem2D dst);
+        void remap_gpu_1c(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
+        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);
 
@@ -73,14 +74,21 @@ namespace cv { namespace gpu
 }}
 
 void cv::gpu::remap(const GpuMat& src, const GpuMat& xmap, const GpuMat& ymap, GpuMat& dst)
-{ 
+{
-    CV_DbgAssert(xmap.data && xmap.cols == ymap.cols && xmap.rows == ymap.rows);
+    typedef void (*remap_gpu_t)(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst);
-    CV_Assert(xmap.type() == CV_32F && ymap.type() == CV_32F);
+    static const remap_gpu_t callers[] = {impl::remap_gpu_1c, 0, impl::remap_gpu_3c};
 
-    dst.create(xmap.size(), src.type());
+    CV_Assert((src.type() == CV_8U || src.type() == CV_8UC3) && xmap.type() == CV_32F && ymap.type() == CV_32F);
-    CV_Assert(dst.data != src.data);   
+
+    GpuMat out;
+    if (dst.data != src.data)
+        out = dst;
+
+    out.create(xmap.size(), src.type());
     
-    impl::remap_gpu(src, xmap, ymap, dst);
+    callers[src.channels() - 1](src, xmap, ymap, out);
+    
+    dst = out;
 }
 
 
@@ -114,7 +122,7 @@ namespace
     void drawColorDisp_caller(const GpuMat& src, GpuMat& dst, int ndisp, const cudaStream_t& stream)
     {        
         GpuMat out;
-        if (&dst != &src)
+        if (dst.data != src.data)
             out = dst;
         out.create(src.size(), CV_8UC4);