prefilter_xsobel option added to stereobm_gpu

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 7ce5e79..d364043 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -326,13 +326,15 @@ namespace cv
         class CV_EXPORTS StereoBM_GPU\r
         {\r
         public:\r
-            enum { BASIC_PRESET=0, PREFILTER_XSOBEL = 1 };\r
+            enum { BASIC_PRESET = 0, PREFILTER_XSOBEL = 1 };\r
+\r
+            enum { DEFAULT_NDISP = 64, DEFAULT_WINSZ = 19 };\r
 \r
             //! the default constructor\r
             StereoBM_GPU();\r
             //! the full constructor taking the camera-specific preset, number of disparities and the SAD window size\r
             //! ndisparities should be multiple of 8. SSD WindowsSize is fixed to 19 now\r
-            StereoBM_GPU(int preset, int ndisparities=0);\r
+            StereoBM_GPU(int preset, int ndisparities = DEFAULT_NDISP, int winSize = DEFAULT_WINSZ);\r
             //! the stereo correspondence operator. Finds the disparity for the specified rectified stereo pair\r
             //! Output disparity has CV_8U type.\r
             void operator() ( const GpuMat& left, const GpuMat& right, GpuMat& disparity);\r
@@ -344,10 +346,13 @@ namespace cv
             // if current GPU will be faster then CPU in this algorithm.\r
             // It queries current active device.\r
             static bool checkIfGpuCallReasonable();\r
-        private:\r
-            GpuMat minSSD;\r
-            int preset;\r
+\r
             int ndisp;\r
+            int winSize;\r
+            int preset;\r
+        private:\r
+            GpuMat minSSD, leBuf, riBuf;\r
+                        \r
         };\r
     }\r
 }\r

diff --git a/modules/gpu/src/cuda/cuda_shared.hpp b/modules/gpu/src/cuda/cuda_shared.hpp
index 23c4c00..272e4f5 100644 (file)
--- a/modules/gpu/src/cuda/cuda_shared.hpp
+++ b/modules/gpu/src/cuda/cuda_shared.hpp
@@ -59,8 +59,7 @@ namespace cv
         namespace impl\r
         {\r
             static inline int divUp(int a, int b) { return (a % b == 0) ? a/b : a/b + 1; }\r
-\r
-            extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, DevMem2D& disp, int maxdisp, DevMem2D_<uint>& minSSD_buf);\r
+           \r
 \r
             extern "C" void set_to_without_mask (const DevMem2D& mat, const double * scalar, int depth, int channels);\r
             extern "C" void set_to_with_mask    (const DevMem2D& mat, const double * scalar, const DevMem2D& mask, int depth, int channels);\r

diff --git a/modules/gpu/src/cuda/stereobm.cu b/modules/gpu/src/cuda/stereobm.cu
index 39d72af..e2427da 100644 (file)
--- a/modules/gpu/src/cuda/stereobm.cu
+++ b/modules/gpu/src/cuda/stereobm.cu
@@ -55,14 +55,14 @@
 #define COL_SSD_SIZE (BLOCK_W + N_DIRTY_PIXELS)\r
 #define SHARED_MEM_SIZE (COL_SSD_SIZE) // amount of shared memory used\r
 \r
+namespace stereobm_gpu \r
+{\r
+\r
 __constant__ unsigned int* cminSSDImage;\r
 __constant__ size_t cminSSD_step;\r
 __constant__ int cwidth;\r
 __constant__ int cheight;\r
 \r
-namespace device_code \r
-{\r
-\r
 __device__ int SQ(int a)\r
 {\r
     return a * a;    \r
@@ -290,29 +290,79 @@ extern "C" __global__ void stereoKernel(unsigned char *left, unsigned char *righ
 \r
 }\r
 \r
-extern "C" void cv::gpu::impl::stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, DevMem2D& disp, int maxdisp, DevMem2D_<unsigned int>& minSSD_buf)\r
-{   \r
-    //cudaSafeCall( cudaFuncSetCacheConfig(&stereoKernel, cudaFuncCachePreferL1) );\r
-    //cudaSafeCall( cudaFuncSetCacheConfig(&stereoKernel, cudaFuncCachePreferShared) );\r
-    \r
-    size_t smem_size = (BLOCK_W + N_DISPARITIES * SHARED_MEM_SIZE) * sizeof(unsigned int);      \r
+namespace cv { namespace gpu { namespace impl\r
+{\r
+    extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int maxdisp, int winsz, const DevMem2D_<unsigned int>& minSSD_buf)\r
+    {   \r
+        //cudaSafeCall( cudaFuncSetCacheConfig(&stereoKernel, cudaFuncCachePreferL1) );\r
+        //cudaSafeCall( cudaFuncSetCacheConfig(&stereoKernel, cudaFuncCachePreferShared) );\r
 \r
-    cudaSafeCall( cudaMemset2D(disp.ptr, disp.step, 0, disp.cols, disp. rows) );\r
-    cudaSafeCall( cudaMemset2D(minSSD_buf.ptr, minSSD_buf.step, 0xFF, minSSD_buf.cols * minSSD_buf.elemSize(), disp. rows) );        \r
+        size_t smem_size = (BLOCK_W + N_DISPARITIES * SHARED_MEM_SIZE) * sizeof(unsigned int);      \r
 \r
-    dim3 grid(1,1,1);\r
-    dim3 threads(BLOCK_W, 1, 1);    \r
-    \r
-    grid.x = divUp(left.cols - maxdisp - 2 * RADIUS, BLOCK_W);\r
-    grid.y = divUp(left.rows - 2 * RADIUS, ROWSperTHREAD);\r
-    \r
-    cudaSafeCall( cudaMemcpyToSymbol(  cwidth, &left.cols, sizeof (left.cols) ) );\r
-    cudaSafeCall( cudaMemcpyToSymbol( cheight, &left.rows, sizeof (left.rows) ) );\r
-    cudaSafeCall( cudaMemcpyToSymbol( cminSSDImage,  &minSSD_buf.ptr, sizeof (minSSD_buf.ptr) ) );\r
-\r
-    size_t minssd_step = minSSD_buf.step/minSSD_buf.elemSize();\r
-    cudaSafeCall( cudaMemcpyToSymbol( cminSSD_step,  &minssd_step, sizeof (minssd_step) ) );\r
-         \r
-    device_code::stereoKernel<<<grid, threads, smem_size>>>(left.ptr, right.ptr, left.step, disp.ptr, disp.step, maxdisp);\r
-    cudaSafeCall( cudaThreadSynchronize() );\r
-}
\ No newline at end of file
+        cudaSafeCall( cudaMemset2D(disp.ptr, disp.step, 0, disp.cols, disp. rows) );\r
+        cudaSafeCall( cudaMemset2D(minSSD_buf.ptr, minSSD_buf.step, 0xFF, minSSD_buf.cols * minSSD_buf.elemSize(), disp. rows) );        \r
+\r
+        dim3 grid(1,1,1);\r
+        dim3 threads(BLOCK_W, 1, 1);    \r
+\r
+        grid.x = divUp(left.cols - maxdisp - 2 * RADIUS, BLOCK_W);\r
+        grid.y = divUp(left.rows - 2 * RADIUS, ROWSperTHREAD);\r
+\r
+        cudaSafeCall( cudaMemcpyToSymbol(  stereobm_gpu::cwidth, &left.cols, sizeof(left.cols) ) );\r
+        cudaSafeCall( cudaMemcpyToSymbol( stereobm_gpu::cheight, &left.rows, sizeof(left.rows) ) );\r
+        cudaSafeCall( cudaMemcpyToSymbol( stereobm_gpu::cminSSDImage, &minSSD_buf.ptr, sizeof(minSSD_buf.ptr) ) );\r
+\r
+        size_t minssd_step = minSSD_buf.step/minSSD_buf.elemSize();\r
+        cudaSafeCall( cudaMemcpyToSymbol( stereobm_gpu::cminSSD_step,  &minssd_step, sizeof(minssd_step) ) );\r
+\r
+        stereobm_gpu::stereoKernel<<<grid, threads, smem_size>>>(left.ptr, right.ptr, left.step, disp.ptr, disp.step, maxdisp);\r
+        cudaSafeCall( cudaThreadSynchronize() );\r
+    }\r
+}}}\r
+\r
+//////////////////////////////////////////////////////////////////////////////////////////////////\r
+/////////////////////////////////////// Sobel Prefiler ///////////////////////////////////////////\r
+//////////////////////////////////////////////////////////////////////////////////////////////////\r
+\r
+namespace stereobm_gpu\r
+{\r
+\r
+texture<unsigned char, 2, cudaReadModeElementType> tex;\r
+\r
+extern "C" __global__ void prefilert_kernel(unsigned char *output, size_t step, int width, int height, int prefilterCap)\r
+{\r
+    int x = blockDim.x * blockIdx.x + threadIdx.x;\r
+    int y = blockDim.y * blockIdx.y + threadIdx.y;\r
+\r
+    if (x < width && y < height)\r
+    {\r
+        int conv = (int)tex2D(tex, x - 1, y - 1) * (-1) + (int)tex2D(tex, x + 1, y - 1) * (1) + \r
+                   (int)tex2D(tex, x - 1, y    ) * (-2) + (int)tex2D(tex, x + 1, y    ) * (2) +\r
+                   (int)tex2D(tex, x - 1, y + 1) * (-1) + (int)tex2D(tex, x + 1, y + 1) * (1);\r
+\r
+\r
+        conv = min(min(max(-prefilterCap, conv), prefilterCap) + prefilterCap, 255);\r
+        output[y * step + x] = conv & 0xFF;\r
+    }\r
+}\r
+\r
+}\r
+\r
+namespace cv { namespace gpu  { namespace impl\r
+{\r
+    extern "C" void prefilter_xsobel(const DevMem2D& input, const DevMem2D& output, int prefilterCap)\r
+    {\r
+        cudaChannelFormatDesc desc = cudaCreateChannelDesc<uchar>();\r
+        cudaSafeCall( cudaBindTexture2D( 0, stereobm_gpu::tex, input.ptr, desc, input.cols, input.rows, input.step ) );  \r
+\r
+        dim3 threads(16, 16, 1);\r
+        dim3 grid(1, 1, 1);\r
+\r
+        grid.x = divUp(input.cols, threads.x);\r
+        grid.y = divUp(input.rows, threads.y);            \r
+\r
+        stereobm_gpu::prefilert_kernel<<<grid, threads>>>(output.ptr, output.step, output.cols, output.rows, prefilterCap);\r
+        cudaSafeCall( cudaThreadSynchronize() );\r
+    }\r
+\r
+}}}
\ No newline at end of file

diff --git a/modules/gpu/src/stereobm_gpu.cpp b/modules/gpu/src/stereobm_gpu.cpp
index ae96700..c557d4e 100644 (file)
--- a/modules/gpu/src/stereobm_gpu.cpp
+++ b/modules/gpu/src/stereobm_gpu.cpp
@@ -56,12 +56,21 @@ void cv::gpu::StereoBM_GPU::operator() ( const GpuMat& left, const GpuMat& right
 \r
 \r
 #else /* !defined (HAVE_CUDA) */\r
+\r
+namespace cv { namespace gpu \r
+{  \r
+    namespace impl \r
+    {\r
+        extern "C" void stereoBM_GPU(const DevMem2D& left, const DevMem2D& right, const DevMem2D& disp, int maxdisp, int winsz, const DevMem2D_<uint>& minSSD_buf);\r
+        extern "C" void prefilter_xsobel(const DevMem2D& input, const DevMem2D& output, int prefilterCap = 31);\r
+    }\r
+}}\r
    \r
-cv::gpu::StereoBM_GPU::StereoBM_GPU() : preset(BASIC_PRESET), ndisp(64)  {}\r
-cv::gpu::StereoBM_GPU::StereoBM_GPU(int preset_, int ndisparities_) : preset(preset_), ndisp(ndisparities_) \r
+cv::gpu::StereoBM_GPU::StereoBM_GPU() : preset(BASIC_PRESET), ndisp(DEFAULT_NDISP), winSize(DEFAULT_WINSZ)  {}\r
+cv::gpu::StereoBM_GPU::StereoBM_GPU(int preset_, int ndisparities_, int winSize_) : preset(preset_), ndisp(ndisparities_), winSize(winSize_) \r
 {\r
-    const int max_supported_ndisp = 1 << (sizeof(unsigned char) * 8);\r
-    CV_Assert(ndisp <= max_supported_ndisp);\r
+    const int max_supported_ndisp = 1 << (sizeof(unsigned char) * 8);    \r
+    CV_Assert(0 < ndisp && ndisp <= max_supported_ndisp);\r
     CV_Assert(ndisp % 8 == 0);\r
 }\r
 \r
@@ -91,14 +100,21 @@ void cv::gpu::StereoBM_GPU::operator() ( const GpuMat& left, const GpuMat& right
     disparity.create(left.size(), CV_8U);\r
     minSSD.create(left.size(), CV_32S);\r
 \r
+    GpuMat le_for_bm =  left;\r
+    GpuMat ri_for_bm = right;\r
+        \r
     if (preset == PREFILTER_XSOBEL)\r
     {\r
-         CV_Assert(!"Not implemented");\r
-    }   \r
-\r
-    DevMem2D disp = disparity;\r
-    DevMem2D_<unsigned int> mssd = minSSD;    \r
-    impl::stereoBM_GPU(left, right, disp, ndisp, mssd);     \r
+        leBuf.create( left.size(),  left.type());\r
+        riBuf.create(right.size(), right.type());\r
+            \r
+        impl::prefilter_xsobel( left, leBuf);\r
+        impl::prefilter_xsobel(right, riBuf);        \r
+\r
+        le_for_bm = leBuf;\r
+        ri_for_bm = riBuf;\r
+    }  \r
+    impl::stereoBM_GPU(le_for_bm, ri_for_bm, disparity, ndisp, winSize, minSSD);     \r
 }\r
 \r
 void cv::gpu::StereoBM_GPU::operator() ( const GpuMat& left, const GpuMat& right, GpuMat& disparity, const CudaStream& stream)\r