#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
\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
\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
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