void cv::gpu::drawColorDisp(const GpuMat&, GpuMat&, int, Stream&) { throw_nogpu(); }\r
void cv::gpu::reprojectImageTo3D(const GpuMat&, GpuMat&, const Mat&, Stream&) { throw_nogpu(); }\r
void cv::gpu::resize(const GpuMat&, GpuMat&, Size, double, double, int, Stream&) { throw_nogpu(); }\r
-void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, const Scalar&, Stream&) { throw_nogpu(); }\r
+void cv::gpu::copyMakeBorder(const GpuMat&, GpuMat&, int, int, int, int, int, const Scalar&, Stream&) { throw_nogpu(); }\r
void cv::gpu::warpAffine(const GpuMat&, GpuMat&, const Mat&, Size, int, Stream&) { throw_nogpu(); }\r
void cv::gpu::warpPerspective(const GpuMat&, GpuMat&, const Mat&, Size, int, Stream&) { throw_nogpu(); }\r
-void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
+void cv::gpu::buildWarpPlaneMaps(Size, Rect, const Mat&, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
void cv::gpu::buildWarpCylindricalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
void cv::gpu::buildWarpSphericalMaps(Size, Rect, const Mat&, const Mat&, float, GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
void cv::gpu::rotate(const GpuMat&, GpuMat&, Size, double, double, double, int, Stream&) { throw_nogpu(); }\r
////////////////////////////////////////////////////////////////////////\r
// copyMakeBorder\r
\r
-void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, const Scalar& value, Stream& s)\r
+namespace cv { namespace gpu { namespace imgproc\r
{\r
- CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1);\r
+ template <typename T, int cn> void copyMakeBorder_gpu(const DevMem2D& src, const DevMem2D& dst, int top, int left, int borderMode, const T* borderValue, cudaStream_t stream);\r
+}}}\r
\r
- dst.create(src.rows + top + bottom, src.cols + left + right, src.type());\r
+namespace\r
+{\r
+ template <typename T, int cn> void copyMakeBorder_caller(const DevMem2D& src, const DevMem2D& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream)\r
+ {\r
+ Scalar_<T> val(saturate_cast<T>(value[0]), saturate_cast<T>(value[1]), saturate_cast<T>(value[2]), saturate_cast<T>(value[3]));\r
\r
- NppiSize srcsz;\r
- srcsz.width = src.cols;\r
- srcsz.height = src.rows;\r
- NppiSize dstsz;\r
- dstsz.width = dst.cols;\r
- dstsz.height = dst.rows;\r
+ imgproc::copyMakeBorder_gpu<T, cn>(src, dst, top, left, borderType, val.val, stream);\r
+ }\r
+}\r
\r
- cudaStream_t stream = StreamAccessor::getStream(s);\r
+void cv::gpu::copyMakeBorder(const GpuMat& src, GpuMat& dst, int top, int bottom, int left, int right, int borderType, const Scalar& value, Stream& s)\r
+{\r
+ CV_Assert(src.depth() <= CV_32F && src.channels() <= 4);\r
+ CV_Assert(borderType == BORDER_REFLECT101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP);\r
\r
- NppStreamHandler h(stream);\r
+ dst.create(src.rows + top + bottom, src.cols + left + right, src.type());\r
+\r
+ cudaStream_t stream = StreamAccessor::getStream(s);\r
\r
- switch (src.type())\r
+ if (borderType == BORDER_CONSTANT && (src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1))\r
{\r
- case CV_8UC1:\r
- {\r
- Npp8u nVal = static_cast<Npp8u>(value[0]);\r
- nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr<Npp8u>(), static_cast<int>(src.step), srcsz,\r
- dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
- break;\r
- }\r
- case CV_8UC4:\r
- {\r
- Npp8u nVal[] = {static_cast<Npp8u>(value[0]), static_cast<Npp8u>(value[1]), static_cast<Npp8u>(value[2]), static_cast<Npp8u>(value[3])};\r
- nppSafeCall( nppiCopyConstBorder_8u_C4R(src.ptr<Npp8u>(), static_cast<int>(src.step), srcsz,\r
- dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
- break;\r
- }\r
- case CV_32SC1:\r
- {\r
- Npp32s nVal = static_cast<Npp32s>(value[0]);\r
- nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr<Npp32s>(), static_cast<int>(src.step), srcsz,\r
- dst.ptr<Npp32s>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
- break;\r
- }\r
- case CV_32FC1:\r
+ NppiSize srcsz;\r
+ srcsz.width = src.cols;\r
+ srcsz.height = src.rows;\r
+\r
+ NppiSize dstsz;\r
+ dstsz.width = dst.cols;\r
+ dstsz.height = dst.rows;\r
+\r
+ NppStreamHandler h(stream);\r
+\r
+ switch (src.type())\r
{\r
- Npp32f val = static_cast<Npp32f>(value[0]);\r
- Npp32s nVal = *(reinterpret_cast<Npp32s*>(&val));\r
- nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr<Npp32s>(), static_cast<int>(src.step), srcsz,\r
- dst.ptr<Npp32s>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
- break;\r
+ case CV_8UC1:\r
+ {\r
+ Npp8u nVal = saturate_cast<Npp8u>(value[0]);\r
+ nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr<Npp8u>(), static_cast<int>(src.step), srcsz,\r
+ dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
+ break;\r
+ }\r
+ case CV_8UC4:\r
+ {\r
+ Npp8u nVal[] = {saturate_cast<Npp8u>(value[0]), saturate_cast<Npp8u>(value[1]), saturate_cast<Npp8u>(value[2]), saturate_cast<Npp8u>(value[3])};\r
+ nppSafeCall( nppiCopyConstBorder_8u_C4R(src.ptr<Npp8u>(), static_cast<int>(src.step), srcsz,\r
+ dst.ptr<Npp8u>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
+ break;\r
+ }\r
+ case CV_32SC1:\r
+ {\r
+ Npp32s nVal = saturate_cast<Npp32s>(value[0]);\r
+ nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr<Npp32s>(), static_cast<int>(src.step), srcsz,\r
+ dst.ptr<Npp32s>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
+ break;\r
+ }\r
+ case CV_32FC1:\r
+ {\r
+ Npp32f val = saturate_cast<Npp32f>(value[0]);\r
+ Npp32s nVal = *(reinterpret_cast<Npp32s*>(&val));\r
+ nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr<Npp32s>(), static_cast<int>(src.step), srcsz,\r
+ dst.ptr<Npp32s>(), static_cast<int>(dst.step), dstsz, top, left, nVal) );\r
+ break;\r
+ }\r
}\r
- default:\r
- CV_Assert(!"Unsupported source type");\r
+\r
+ if (stream == 0)\r
+ cudaSafeCall( cudaDeviceSynchronize() );\r
}\r
+ else\r
+ {\r
+ typedef void (*caller_t)(const DevMem2D& src, const DevMem2D& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream);\r
+ static const caller_t callers[6][4] = \r
+ {\r
+ { copyMakeBorder_caller<uchar, 1> , 0/*copyMakeBorder_caller<uchar, 2>*/ , copyMakeBorder_caller<uchar, 3> , copyMakeBorder_caller<uchar, 4>},\r
+ {0/*copyMakeBorder_caller<schar, 1>*/, 0/*copyMakeBorder_caller<schar, 2>*/ , 0/*copyMakeBorder_caller<schar, 3>*/, 0/*copyMakeBorder_caller<schar, 4>*/},\r
+ { copyMakeBorder_caller<ushort, 1> , 0/*copyMakeBorder_caller<ushort, 2>*/, copyMakeBorder_caller<ushort, 3> , copyMakeBorder_caller<ushort, 4>},\r
+ { copyMakeBorder_caller<short, 1> , 0/*copyMakeBorder_caller<short, 2>*/ , copyMakeBorder_caller<short, 3> , copyMakeBorder_caller<short, 4>},\r
+ {0/*copyMakeBorder_caller<int, 1>*/ , 0/*copyMakeBorder_caller<int, 2>*/ , 0/*copyMakeBorder_caller<int, 3>*/ , 0/*copyMakeBorder_caller<int, 4>*/},\r
+ { copyMakeBorder_caller<float, 1> , 0/*copyMakeBorder_caller<float, 2>*/ , copyMakeBorder_caller<float, 3> , copyMakeBorder_caller<float ,4>}\r
+ };\r
\r
- if (stream == 0)\r
- cudaSafeCall( cudaDeviceSynchronize() );\r
+ caller_t func = callers[src.depth()][src.channels() - 1];\r
+ CV_Assert(func != 0);\r
+\r
+ int gpuBorderType;\r
+ CV_Assert(tryConvertToGpuBorderType(borderType, gpuBorderType));\r
+\r
+ func(src, dst, top, left, gpuBorderType, value, stream);\r
+ }\r
}\r
\r
////////////////////////////////////////////////////////////////////////\r
namespace cv { namespace gpu { namespace imgproc\r
{\r
void buildWarpPlaneMaps(int tl_u, int tl_v, DevMem2Df map_x, DevMem2Df map_y,\r
- const float k_rinv[9], const float r_kinv[9], float scale,\r
+ const float k_rinv[9], const float r_kinv[9], const float t[3], float scale,\r
cudaStream_t stream);\r
}}}\r
\r
-void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, float scale,\r
- GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
+void cv::gpu::buildWarpPlaneMaps(Size src_size, Rect dst_roi, const Mat &K, const Mat& R, const Mat &T, \r
+ float scale, GpuMat& map_x, GpuMat& map_y, Stream& stream)\r
{\r
CV_Assert(K.size() == Size(3,3) && K.type() == CV_32F);\r
CV_Assert(R.size() == Size(3,3) && R.type() == CV_32F);\r
+ CV_Assert((T.size() == Size(3,1) || T.size() == Size(1,3)) && T.type() == CV_32F && T.isContinuous());\r
\r
Mat K_Rinv = K * R.t();\r
Mat R_Kinv = R * K.inv();\r
\r
map_x.create(dst_roi.size(), CV_32F);\r
map_y.create(dst_roi.size(), CV_32F);\r
- imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(),\r
- scale, StreamAccessor::getStream(stream));\r
+ imgproc::buildWarpPlaneMaps(dst_roi.tl().x, dst_roi.tl().y, map_x, map_y, K_Rinv.ptr<float>(), R_Kinv.ptr<float>(), \r
+ T.ptr<float>(), scale, StreamAccessor::getStream(stream));\r
}\r
\r
//////////////////////////////////////////////////////////////////////////////\r
projects / profile / ivi / opencv.git / blobdiff