diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 2081ba3f24..25ca7d2d89 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/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, diff --git a/modules/gpu/src/beliefpropagation_gpu.cpp b/modules/gpu/src/beliefpropagation_gpu.cpp index d5e3fab65a..31fe7c19a5 100644 --- a/modules/gpu/src/beliefpropagation_gpu.cpp +++ b/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(mm)) + 1) * 4 / 5; diff --git a/modules/gpu/src/constantspacebp_gpu.cpp b/modules/gpu/src/constantspacebp_gpu.cpp index d14f7aea7a..30b5c5d7bd 100644 --- a/modules/gpu/src/constantspacebp_gpu.cpp +++ b/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++; diff --git a/modules/gpu/src/cuda/imgproc.cu b/modules/gpu/src/cuda/imgproc.cu index a08e34ab95..5638ccd69b 100644 --- a/modules/gpu/src/cuda/imgproc.cu +++ b/modules/gpu/src/cuda/imgproc.cu @@ -44,13 +44,12 @@ using namespace cv::gpu; - /////////////////////////////////// Remap /////////////////////////////////////////////// namespace imgproc { texture 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_& xmap, const DevMem2D_& 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.y = divUp(dst.rows, block.y); + grid.x = divUp(dst.cols, threads.x); + 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(); - 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<<>>(xmap.ptr, ymap.ptr, xmap.step, dst.ptr, dst.step, dst.cols, dst.rows); + imgproc::remap_1c<<>>(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<<>>(src.ptr, src.step, xmap.ptr, ymap.ptr, xmap.step, dst.ptr, dst.step, dst.cols, dst.rows); + + cudaSafeCall( cudaThreadSynchronize() ); + } }}} diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp index ee06fe5939..ab7efa6a60 100644 --- a/modules/gpu/src/imgproc_gpu.cpp +++ b/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_& xmap, const DevMem2D_& 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); - CV_Assert(xmap.type() == CV_32F && ymap.type() == CV_32F); +{ + typedef void (*remap_gpu_t)(const DevMem2D& src, const DevMem2Df& xmap, const DevMem2Df& ymap, DevMem2D dst); + static const remap_gpu_t callers[] = {impl::remap_gpu_1c, 0, impl::remap_gpu_3c}; - dst.create(xmap.size(), src.type()); - CV_Assert(dst.data != src.data); + CV_Assert((src.type() == CV_8U || src.type() == CV_8UC3) && xmap.type() == CV_32F && ymap.type() == CV_32F); + + 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);