From 32d578f5f043306057928032b9d5713b7e2535a9 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 29 Jul 2013 14:45:06 +0400 Subject: [PATCH 01/35] fixed gridTransform overloads problems --- modules/core/src/cuda/gpu_mat.cu | 10 +- .../opencv2/cudev/grid/detail/transform.hpp | 10 +- .../include/opencv2/cudev/grid/transform.hpp | 136 +++++++++--------- 3 files changed, 78 insertions(+), 78 deletions(-) diff --git a/modules/core/src/cuda/gpu_mat.cu b/modules/core/src/cuda/gpu_mat.cu index d9332b7332..af0a4f5d0d 100644 --- a/modules/core/src/cuda/gpu_mat.cu +++ b/modules/core/src/cuda/gpu_mat.cu @@ -216,7 +216,7 @@ namespace template void copyWithMask(const GpuMat& src, const GpuMat& dst, const GpuMat& mask, Stream& stream) { - gridTransform_< CopyToPolicy::elem_type)> >(globPtr(src), globPtr(dst), identity(), globPtr(mask), stream); + gridTransformUnary_< CopyToPolicy::elem_type)> >(globPtr(src), globPtr(dst), identity(), globPtr(mask), stream); } } @@ -268,14 +268,14 @@ namespace void setToWithOutMask(const GpuMat& mat, Scalar _scalar, Stream& stream) { Scalar_::elem_type> scalar = _scalar; - gridTransform(constantPtr(VecTraits::make(scalar.val), mat.rows, mat.cols), globPtr(mat), identity(), stream); + gridTransformUnary(constantPtr(VecTraits::make(scalar.val), mat.rows, mat.cols), globPtr(mat), identity(), stream); } template void setToWithMask(const GpuMat& mat, const GpuMat& mask, Scalar _scalar, Stream& stream) { Scalar_::elem_type> scalar = _scalar; - gridTransform(constantPtr(VecTraits::make(scalar.val), mat.rows, mat.cols), globPtr(mat), identity(), globPtr(mask), stream); + gridTransformUnary(constantPtr(VecTraits::make(scalar.val), mat.rows, mat.cols), globPtr(mat), identity(), globPtr(mask), stream); } } @@ -382,7 +382,7 @@ namespace typedef typename LargerType::type larger_elem_type; typedef typename LargerType::type scalar_type; - gridTransform_< ConvertToPolicy >(globPtr(src), globPtr(dst), saturate_cast_func(), stream); + gridTransformUnary_< ConvertToPolicy >(globPtr(src), globPtr(dst), saturate_cast_func(), stream); } template struct Convertor : unary_function @@ -408,7 +408,7 @@ namespace op.alpha = cv::saturate_cast(alpha); op.beta = cv::saturate_cast(beta); - gridTransform_< ConvertToPolicy >(globPtr(src), globPtr(dst), op, stream); + gridTransformUnary_< ConvertToPolicy >(globPtr(src), globPtr(dst), op, stream); } } diff --git a/modules/cudev/include/opencv2/cudev/grid/detail/transform.hpp b/modules/cudev/include/opencv2/cudev/grid/detail/transform.hpp index 0e0df0794a..d13585398b 100644 --- a/modules/cudev/include/opencv2/cudev/grid/detail/transform.hpp +++ b/modules/cudev/include/opencv2/cudev/grid/detail/transform.hpp @@ -217,7 +217,7 @@ namespace grid_transform_detail } template - __global__ void transformSmart(const GlobPtr src1_, const GlobPtr src2_, PtrStep dst_, const BinOp op, const MaskPtr mask, const int rows, const int cols) + __global__ void transformSmart(const GlobPtr src1_, const GlobPtr src2_, GlobPtr dst_, const BinOp op, const MaskPtr mask, const int rows, const int cols) { typedef typename MakeVec::type read_type1; typedef typename MakeVec::type read_type2; @@ -345,25 +345,25 @@ namespace grid_transform_detail }; template - __host__ void transform(const SrcPtr& src, const GlobPtr& dst, const UnOp& op, const MaskPtr& mask, int rows, int cols, cudaStream_t stream) + __host__ void transform_unary(const SrcPtr& src, const GlobPtr& dst, const UnOp& op, const MaskPtr& mask, int rows, int cols, cudaStream_t stream) { TransformDispatcher::call(src, dst, op, mask, rows, cols, stream); } template - __host__ void transform(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtr& dst, const BinOp& op, const MaskPtr& mask, int rows, int cols, cudaStream_t stream) + __host__ void transform_binary(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtr& dst, const BinOp& op, const MaskPtr& mask, int rows, int cols, cudaStream_t stream) { TransformDispatcher::call(src1, src2, dst, op, mask, rows, cols, stream); } template - __host__ void transform(const GlobPtr& src, const GlobPtr& dst, const UnOp& op, const MaskPtr& mask, int rows, int cols, cudaStream_t stream) + __host__ void transform_unary(const GlobPtr& src, const GlobPtr& dst, const UnOp& op, const MaskPtr& mask, int rows, int cols, cudaStream_t stream) { TransformDispatcher::cn == 1 && VecTraits::cn == 1 && Policy::shift != 1, Policy>::call(src, dst, op, mask, rows, cols, stream); } template - __host__ void transform(const GlobPtr& src1, const GlobPtr& src2, const GlobPtr& dst, const BinOp& op, const MaskPtr& mask, int rows, int cols, cudaStream_t stream) + __host__ void transform_binary(const GlobPtr& src1, const GlobPtr& src2, const GlobPtr& dst, const BinOp& op, const MaskPtr& mask, int rows, int cols, cudaStream_t stream) { TransformDispatcher::cn == 1 && VecTraits::cn == 1 && VecTraits::cn == 1 && Policy::shift != 1, Policy>::call(src1, src2, dst, op, mask, rows, cols, stream); } diff --git a/modules/cudev/include/opencv2/cudev/grid/transform.hpp b/modules/cudev/include/opencv2/cudev/grid/transform.hpp index 0da5e2d8a4..62555ab5ae 100644 --- a/modules/cudev/include/opencv2/cudev/grid/transform.hpp +++ b/modules/cudev/include/opencv2/cudev/grid/transform.hpp @@ -58,7 +58,7 @@ namespace cv { namespace cudev { template -__host__ void gridTransform_(const SrcPtr& src, GpuMat_& dst, const UnOp& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformUnary_(const SrcPtr& src, GpuMat_& dst, const UnOp& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { const int rows = getRows(src); const int cols = getCols(src); @@ -67,11 +67,11 @@ __host__ void gridTransform_(const SrcPtr& src, GpuMat_& dst, const UnO dst.create(rows, cols); - grid_transform_detail::transform(shrinkPtr(src), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream)); + grid_transform_detail::transform_unary(shrinkPtr(src), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream)); } template -__host__ void gridTransform_(const SrcPtr& src, const GlobPtrSz& dst, const UnOp& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformUnary_(const SrcPtr& src, const GlobPtrSz& dst, const UnOp& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { const int rows = getRows(src); const int cols = getCols(src); @@ -79,33 +79,33 @@ __host__ void gridTransform_(const SrcPtr& src, const GlobPtrSz& dst, c CV_Assert( getRows(dst) == rows && getCols(dst) == cols ); CV_Assert( getRows(mask) == rows && getCols(mask) == cols ); - grid_transform_detail::transform(shrinkPtr(src), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream)); + grid_transform_detail::transform_unary(shrinkPtr(src), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream)); } template -__host__ void gridTransform_(const SrcPtr& src, GpuMat_& dst, const UnOp& op, Stream& stream = Stream::Null()) +__host__ void gridTransformUnary_(const SrcPtr& src, GpuMat_& dst, const UnOp& op, Stream& stream = Stream::Null()) { const int rows = getRows(src); const int cols = getCols(src); dst.create(rows, cols); - grid_transform_detail::transform(shrinkPtr(src), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream)); + grid_transform_detail::transform_unary(shrinkPtr(src), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream)); } template -__host__ void gridTransform_(const SrcPtr& src, const GlobPtrSz& dst, const UnOp& op, Stream& stream = Stream::Null()) +__host__ void gridTransformUnary_(const SrcPtr& src, const GlobPtrSz& dst, const UnOp& op, Stream& stream = Stream::Null()) { const int rows = getRows(src); const int cols = getCols(src); CV_Assert( getRows(dst) == rows && getCols(dst) == cols ); - grid_transform_detail::transform(shrinkPtr(src), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream)); + grid_transform_detail::transform_unary(shrinkPtr(src), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream)); } template -__host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_& dst, const BinOp& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_& dst, const BinOp& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { const int rows = getRows(src1); const int cols = getCols(src1); @@ -115,11 +115,11 @@ __host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream)); + grid_transform_detail::transform_binary(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream)); } template -__host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz& dst, const BinOp& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz& dst, const BinOp& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { const int rows = getRows(src1); const int cols = getCols(src1); @@ -128,11 +128,11 @@ __host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, const Glo CV_Assert( getRows(src2) == rows && getCols(src2) == cols ); CV_Assert( getRows(mask) == rows && getCols(mask) == cols ); - grid_transform_detail::transform(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream)); + grid_transform_detail::transform_binary(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, shrinkPtr(mask), rows, cols, StreamAccessor::getStream(stream)); } template -__host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_& dst, const BinOp& op, Stream& stream = Stream::Null()) +__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_& dst, const BinOp& op, Stream& stream = Stream::Null()) { const int rows = getRows(src1); const int cols = getCols(src1); @@ -141,11 +141,11 @@ __host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream)); + grid_transform_detail::transform_binary(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream)); } template -__host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GlobPtrSz& dst, const BinOp& op, Stream& stream = Stream::Null()) +__host__ void gridTransformBinary_(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz& dst, const BinOp& op, Stream& stream = Stream::Null()) { const int rows = getRows(src1); const int cols = getCols(src1); @@ -153,11 +153,11 @@ __host__ void gridTransform_(const SrcPtr1& src1, const SrcPtr2& src2, GlobPtrSz CV_Assert( getRows(dst) == rows && getCols(dst) == cols ); CV_Assert( getRows(src2) == rows && getCols(src2) == cols ); - grid_transform_detail::transform(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream)); + grid_transform_detail::transform_binary(shrinkPtr(src1), shrinkPtr(src2), shrinkPtr(dst), op, WithOutMask(), rows, cols, StreamAccessor::getStream(stream)); } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 2, "" ); @@ -178,7 +178,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMa } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 2, "" ); @@ -198,7 +198,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, Glob } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 2, "" ); @@ -217,7 +217,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMa } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 2, "" ); @@ -236,7 +236,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, Glob } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 3, "" ); @@ -258,7 +258,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMa } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 3, "" ); @@ -279,7 +279,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, Glob } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 3, "" ); @@ -299,7 +299,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMa } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 3, "" ); @@ -319,7 +319,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, Glob } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 4, "" ); @@ -342,7 +342,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMa } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 4, "" ); @@ -364,7 +364,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, Glob } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 4, "" ); @@ -385,7 +385,7 @@ __host__ void gridTransform_(const SrcPtr& src, const tuple< GpuMat_&, GpuMa } template -__host__ void gridTransform_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple_(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { CV_StaticAssert( tuple_size::value == 4, "" ); @@ -417,123 +417,123 @@ struct DefaultTransformPolicy }; template -__host__ void gridTransform(const SrcPtr& src, GpuMat_& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformUnary(const SrcPtr& src, GpuMat_& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, mask, stream); + gridTransformUnary_(src, dst, op, mask, stream); } template -__host__ void gridTransform(const SrcPtr& src, const GlobPtrSz& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformUnary(const SrcPtr& src, const GlobPtrSz& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, mask, stream); + gridTransformUnary_(src, dst, op, mask, stream); } template -__host__ void gridTransform(const SrcPtr& src, GpuMat_& dst, const Op& op, Stream& stream = Stream::Null()) +__host__ void gridTransformUnary(const SrcPtr& src, GpuMat_& dst, const Op& op, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, stream); + gridTransformUnary_(src, dst, op, stream); } template -__host__ void gridTransform(const SrcPtr& src, const GlobPtrSz& dst, const Op& op, Stream& stream = Stream::Null()) +__host__ void gridTransformUnary(const SrcPtr& src, const GlobPtrSz& dst, const Op& op, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, stream); + gridTransformUnary_(src, dst, op, stream); } template -__host__ void gridTransform(const SrcPtr1& src1, const SrcPtr1& src2, GpuMat_& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { - gridTransform_(src1, src2, dst, op, mask, stream); + gridTransformBinary_(src1, src2, dst, op, mask, stream); } template -__host__ void gridTransform(const SrcPtr1& src1, const SrcPtr1& src2, const GlobPtrSz& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz& dst, const Op& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { - gridTransform_(src1, src2, dst, op, mask, stream); + gridTransformBinary_(src1, src2, dst, op, mask, stream); } template -__host__ void gridTransform(const SrcPtr1& src1, const SrcPtr1& src2, GpuMat_& dst, const Op& op, Stream& stream = Stream::Null()) +__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2, GpuMat_& dst, const Op& op, Stream& stream = Stream::Null()) { - gridTransform_(src1, src2, dst, op, stream); + gridTransformBinary_(src1, src2, dst, op, stream); } template -__host__ void gridTransform(const SrcPtr1& src1, const SrcPtr1& src2, const GlobPtrSz& dst, const Op& op, Stream& stream = Stream::Null()) +__host__ void gridTransformBinary(const SrcPtr1& src1, const SrcPtr2& src2, const GlobPtrSz& dst, const Op& op, Stream& stream = Stream::Null()) { - gridTransform_(src1, src2, dst, op, stream); + gridTransformBinary_(src1, src2, dst, op, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, mask, stream); + gridTransformTuple_(src, dst, op, mask, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, mask, stream); + gridTransformTuple_(src, dst, op, mask, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, stream); + gridTransformTuple_(src, dst, op, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, stream); + gridTransformTuple_(src, dst, op, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, mask, stream); + gridTransformTuple_(src, dst, op, mask, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, mask, stream); + gridTransformTuple_(src, dst, op, mask, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, stream); + gridTransformTuple_(src, dst, op, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, stream); + gridTransformTuple_(src, dst, op, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, mask, stream); + gridTransformTuple_(src, dst, op, mask, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, const MaskPtr& mask, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, mask, stream); + gridTransformTuple_(src, dst, op, mask, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GpuMat_&, GpuMat_&, GpuMat_&, GpuMat_& >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, stream); + gridTransformTuple_(src, dst, op, stream); } template -__host__ void gridTransform(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) +__host__ void gridTransformTuple(const SrcPtr& src, const tuple< GlobPtrSz, GlobPtrSz, GlobPtrSz, GlobPtrSz >& dst, const OpTuple& op, Stream& stream = Stream::Null()) { - gridTransform_(src, dst, op, stream); + gridTransformTuple_(src, dst, op, stream); } }} From 9c5da2ea22ebcece572c700d738e8934a86ebbd9 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Fri, 23 Aug 2013 18:28:13 +0400 Subject: [PATCH 02/35] used new device layer for cv::gpu::add --- modules/cudaarithm/CMakeLists.txt | 2 +- modules/cudaarithm/src/cuda/add_mat.cu | 274 ++++++++++-------- modules/cudaarithm/src/cuda/add_scalar.cu | 204 +++++++------ modules/cudaarithm/src/element_operations.cpp | 243 +--------------- .../opencv2/cudev/grid/detail/integral.hpp | 4 +- 5 files changed, 280 insertions(+), 447 deletions(-) diff --git a/modules/cudaarithm/CMakeLists.txt b/modules/cudaarithm/CMakeLists.txt index 67a7ff922f..e676fa6735 100644 --- a/modules/cudaarithm/CMakeLists.txt +++ b/modules/cudaarithm/CMakeLists.txt @@ -6,7 +6,7 @@ set(the_description "CUDA-accelerated Operations on Matrices") ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4324 /wd4512 -Wundef -Wmissing-declarations) -ocv_add_module(cudaarithm opencv_core OPTIONAL opencv_cudalegacy) +ocv_add_module(cudaarithm opencv_core OPTIONAL opencv_cudev opencv_cudalegacy) ocv_module_include_directories() ocv_glob_module_sources() diff --git a/modules/cudaarithm/src/cuda/add_mat.cu b/modules/cudaarithm/src/cuda/add_mat.cu index 1270438208..6e7a7925fd 100644 --- a/modules/cudaarithm/src/cuda/add_mat.cu +++ b/modules/cudaarithm/src/cuda/add_mat.cu @@ -40,146 +40,186 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void addMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& _stream, int); + +namespace { - struct VAdd4 : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, uint b) const - { - return vadd4(a, b); - } - - __host__ __device__ __forceinline__ VAdd4() {} - __host__ __device__ __forceinline__ VAdd4(const VAdd4&) {} - }; - - struct VAdd2 : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, uint b) const - { - return vadd2(a, b); - } - - __host__ __device__ __forceinline__ VAdd2() {} - __host__ __device__ __forceinline__ VAdd2(const VAdd2&) {} - }; - - template struct AddMat : binary_function + template struct AddOp1 : binary_function { __device__ __forceinline__ D operator ()(T a, T b) const { return saturate_cast(a + b); } - - __host__ __device__ __forceinline__ AddMat() {} - __host__ __device__ __forceinline__ AddMat(const AddMat&) {} }; -} - -namespace cv { namespace cuda { namespace device -{ - template <> struct TransformFunctorTraits< arithm::VAdd4 > : arithm::ArithmFuncTraits - { - }; - - template <> struct TransformFunctorTraits< arithm::VAdd2 > : arithm::ArithmFuncTraits - { - }; - - template struct TransformFunctorTraits< arithm::AddMat > : arithm::ArithmFuncTraits - { - }; -}}} - -namespace arithm -{ - void addMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, VAdd4(), WithOutMask(), stream); - } - - void addMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, VAdd2(), WithOutMask(), stream); - } template - void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) + void addMat_v1(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, Stream& stream) { if (mask.data) - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AddMat(), mask, stream); + gridTransformBinary(globPtr(src1), globPtr(src2), globPtr(dst), AddOp1(), globPtr(mask), stream); else - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AddMat(), WithOutMask(), stream); + gridTransformBinary(globPtr(src1), globPtr(src2), globPtr(dst), AddOp1(), stream); } - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + struct AddOp2 : binary_function + { + __device__ __forceinline__ uint operator ()(uint a, uint b) const + { + return vadd2(a, b); + } + }; - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + void addMat_v2(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) + { + const int vcols = src1.cols >> 1; - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + GlobPtrSz src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols); + GlobPtrSz src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols); + GlobPtrSz dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + gridTransformBinary(src1_, src2_, dst_, AddOp2(), stream); + } - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + struct AddOp4 : binary_function + { + __device__ __forceinline__ uint operator ()(uint a, uint b) const + { + return vadd4(a, b); + } + }; - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + void addMat_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) + { + const int vcols = src1.cols >> 2; - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + GlobPtrSz src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols); + GlobPtrSz src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols); + GlobPtrSz dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols); + + gridTransformBinary(src1_, src2_, dst_, AddOp4(), stream); + } } -#endif // CUDA_DISABLER +void addMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int) +{ + typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, Stream& stream); + static const func_t funcs[7][7] = + { + { + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1 + }, + { + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1 + }, + { + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1 + }, + { + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1, + addMat_v1 + }, + { + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + addMat_v1, + addMat_v1, + addMat_v1 + }, + { + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + addMat_v1, + addMat_v1 + }, + { + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + 0 /*addMat_v1*/, + addMat_v1 + } + }; + + const int sdepth = src1.depth(); + const int ddepth = dst.depth(); + + CV_DbgAssert( sdepth < 7 && ddepth < 7 ); + + GpuMat src1_ = src1.reshape(1); + GpuMat src2_ = src2.reshape(1); + GpuMat dst_ = dst.reshape(1); + + if (mask.empty() && (sdepth == CV_8U || sdepth == CV_16U) && ddepth == sdepth) + { + const intptr_t src1ptr = reinterpret_cast(src1_.data); + const intptr_t src2ptr = reinterpret_cast(src2_.data); + const intptr_t dstptr = reinterpret_cast(dst_.data); + + const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0; + + if (isAllAligned) + { + if (sdepth == CV_8U && (src1_.cols & 3) == 0) + { + addMat_v4(src1_, src2_, dst_, stream); + return; + } + else if (sdepth == CV_16U && (src1_.cols & 1) == 0) + { + addMat_v2(src1_, src2_, dst_, stream); + return; + } + } + } + + const func_t func = funcs[sdepth][ddepth]; + + if (!func) + CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); + + func(src1_, src2_, dst_, mask, stream); +} + +#endif diff --git a/modules/cudaarithm/src/cuda/add_scalar.cu b/modules/cudaarithm/src/cuda/add_scalar.cu index 680061be11..e0788e9bdd 100644 --- a/modules/cudaarithm/src/cuda/add_scalar.cu +++ b/modules/cudaarithm/src/cuda/add_scalar.cu @@ -40,109 +40,141 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void addScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int); + +namespace { - template struct AddScalar : unary_function + template struct AddScalarOp : unary_function { - S val; + ScalarType val; - __host__ explicit AddScalar(S val_) : val(val_) {} - - __device__ __forceinline__ D operator ()(T a) const + __device__ __forceinline__ DstType operator ()(SrcType a) const { - return saturate_cast(a + val); + return saturate_cast(saturate_cast(a) + val); } }; -} -namespace cv { namespace cuda { namespace device -{ - template struct TransformFunctorTraits< arithm::AddScalar > : arithm::ArithmFuncTraits + template struct TransformPolicy : DefaultTransformPolicy { }; -}}} - -namespace arithm -{ - template - void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) + template <> struct TransformPolicy : DefaultTransformPolicy { - AddScalar op(static_cast(val)); + enum { + shift = 1 + }; + }; + + template + void addScalarImpl(const GpuMat& src, cv::Scalar value, GpuMat& dst, const GpuMat& mask, Stream& stream) + { + typedef typename MakeVec::cn>::type ScalarType; + + cv::Scalar_ value_ = value; + + AddScalarOp op; + op.val = VecTraits::make(value_.val); if (mask.data) - device::transform((PtrStepSz) src1, (PtrStepSz) dst, op, mask, stream); + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, globPtr(mask), stream); else - device::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, stream); } - - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); } -#endif // CUDA_DISABLER +void addScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int) +{ + typedef void (*func_t)(const GpuMat& src, cv::Scalar val, GpuMat& dst, const GpuMat& mask, Stream& stream); + static const func_t funcs[7][7][4] = + { + { + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl} + }, + { + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl} + }, + { + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl} + }, + { + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl} + }, + { + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl} + }, + { + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl} + }, + { + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/, 0 /*addScalarImpl*/}, + {addScalarImpl, addScalarImpl, addScalarImpl, addScalarImpl} + } + }; + + const int sdepth = src.depth(); + const int ddepth = dst.depth(); + const int cn = src.channels(); + + CV_DbgAssert( sdepth < 7 && ddepth < 7 && cn <= 4 ); + + const func_t func = funcs[sdepth][ddepth][cn - 1]; + + if (!func) + CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); + + func(src, val, dst, mask, stream); +} + +#endif diff --git a/modules/cudaarithm/src/element_operations.cpp b/modules/cudaarithm/src/element_operations.cpp index 715771b671..b3711dcc1d 100644 --- a/modules/cudaarithm/src/element_operations.cpp +++ b/modules/cudaarithm/src/element_operations.cpp @@ -336,248 +336,9 @@ namespace //////////////////////////////////////////////////////////////////////// // add -namespace arithm -{ - void addMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream); - void addMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream); +void addMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& _stream, int); - template - void addMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); -} - -static void addMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& _stream, int) -{ - typedef void (*func_t)(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - static const func_t funcs[7][7] = - { - { - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat - }, - { - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat - }, - { - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat - }, - { - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat, - arithm::addMat - }, - { - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - arithm::addMat, - arithm::addMat, - arithm::addMat - }, - { - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - arithm::addMat, - arithm::addMat - }, - { - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - 0 /*arithm::addMat*/, - arithm::addMat - } - }; - - const int sdepth = src1.depth(); - const int ddepth = dst.depth(); - const int cn = src1.channels(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - PtrStepSzb src1_(src1.rows, src1.cols * cn, src1.data, src1.step); - PtrStepSzb src2_(src1.rows, src1.cols * cn, src2.data, src2.step); - PtrStepSzb dst_(src1.rows, src1.cols * cn, dst.data, dst.step); - - if (mask.empty() && (sdepth == CV_8U || sdepth == CV_16U) && ddepth == sdepth) - { - const intptr_t src1ptr = reinterpret_cast(src1_.data); - const intptr_t src2ptr = reinterpret_cast(src2_.data); - const intptr_t dstptr = reinterpret_cast(dst_.data); - - const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0; - - if (isAllAligned) - { - if (sdepth == CV_8U && (src1_.cols & 3) == 0) - { - const int vcols = src1_.cols >> 2; - - arithm::addMat_v4(PtrStepSz(src1_.rows, vcols, (unsigned int*) src1_.data, src1_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) src2_.data, src2_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) dst_.data, dst_.step), - stream); - - return; - } - else if (sdepth == CV_16U && (src1_.cols & 1) == 0) - { - const int vcols = src1_.cols >> 1; - - arithm::addMat_v2(PtrStepSz(src1_.rows, vcols, (unsigned int*) src1_.data, src1_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) src2_.data, src2_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) dst_.data, dst_.step), - stream); - - return; - } - } - } - - const func_t func = funcs[sdepth][ddepth]; - - if (!func) - CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - - func(src1_, src2_, dst_, mask, stream); -} - -namespace arithm -{ - template - void addScalar(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); -} - -static void addScalar(const GpuMat& src, Scalar val, bool, GpuMat& dst, const GpuMat& mask, double, Stream& _stream, int) -{ - typedef void (*func_t)(PtrStepSzb src1, double val, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - static const func_t funcs[7][7] = - { - { - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar - }, - { - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar - }, - { - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar - }, - { - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar - }, - { - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - arithm::addScalar, - arithm::addScalar, - arithm::addScalar - }, - { - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - arithm::addScalar, - arithm::addScalar - }, - { - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - 0 /*arithm::addScalar*/, - arithm::addScalar - } - }; - - typedef void (*npp_func_t)(const PtrStepSzb src, Scalar sc, PtrStepb dst, cudaStream_t stream); - static const npp_func_t npp_funcs[7][4] = - { - {NppArithmScalar::call, 0 , NppArithmScalar::call, NppArithmScalar::call}, - {0 , 0 , 0 , 0 }, - {NppArithmScalar::call, 0 , NppArithmScalar::call, NppArithmScalar::call}, - {NppArithmScalar::call, NppArithmScalar::call, NppArithmScalar::call, NppArithmScalar::call}, - {NppArithmScalar::call, NppArithmScalar::call, NppArithmScalar::call, 0 }, - {NppArithmScalar::call, NppArithmScalar::call, NppArithmScalar::call, NppArithmScalar::call}, - {0 , 0 , 0 , 0 } - }; - - const int sdepth = src.depth(); - const int ddepth = dst.depth(); - const int cn = src.channels(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - const npp_func_t npp_func = npp_funcs[sdepth][cn - 1]; - if (ddepth == sdepth && cn > 1 && npp_func != 0) - { - npp_func(src, val, dst, stream); - return; - } - - CV_Assert( cn == 1 ); - - const func_t func = funcs[sdepth][ddepth]; - - if (!func) - CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - - func(src, val[0], dst, mask, stream); -} +void addScalar(const GpuMat& src, Scalar val, bool, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int); void cv::cuda::add(InputArray src1, InputArray src2, OutputArray dst, InputArray mask, int dtype, Stream& stream) { diff --git a/modules/cudev/include/opencv2/cudev/grid/detail/integral.hpp b/modules/cudev/include/opencv2/cudev/grid/detail/integral.hpp index b06d7ddbb6..5c90e99893 100644 --- a/modules/cudev/include/opencv2/cudev/grid/detail/integral.hpp +++ b/modules/cudev/include/opencv2/cudev/grid/detail/integral.hpp @@ -594,7 +594,7 @@ namespace integral_detail CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() ); } - __host__ static void integral(const GlobPtr src, GlobPtr dst, int rows, int cols, cudaStream_t stream) + __host__ static void integral(const GlobPtr& src, const GlobPtr& dst, int rows, int cols, cudaStream_t stream) { if (deviceSupports(FEATURE_SET_COMPUTE_30) && (cols % 16 == 0) @@ -614,7 +614,7 @@ namespace integral_detail CV_CUDEV_SAFE_CALL( cudaDeviceSynchronize() ); } - __host__ static void integral(const GlobPtr src, GlobPtr dst, int rows, int cols, cudaStream_t stream) + __host__ __forceinline__ void integral(const GlobPtr& src, const GlobPtr& dst, int rows, int cols, cudaStream_t stream) { GlobPtr dstui = globPtr((uint*) dst.data, dst.step); integral(src, dstui, rows, cols, stream); From 156f86ea0b66363ba470c09f4f564e76b18796a2 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 29 Jul 2013 15:17:38 +0400 Subject: [PATCH 03/35] used new device layer for cv::gpu::subtract --- modules/cudaarithm/src/cuda/sub_mat.cu | 274 ++++++++++-------- modules/cudaarithm/src/cuda/sub_scalar.cu | 230 +++++++++------ modules/cudaarithm/src/element_operations.cpp | 243 +--------------- 3 files changed, 301 insertions(+), 446 deletions(-) diff --git a/modules/cudaarithm/src/cuda/sub_mat.cu b/modules/cudaarithm/src/cuda/sub_mat.cu index 873b73c46a..ec8d229229 100644 --- a/modules/cudaarithm/src/cuda/sub_mat.cu +++ b/modules/cudaarithm/src/cuda/sub_mat.cu @@ -40,146 +40,186 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void subMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& _stream, int); + +namespace { - struct VSub4 : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, uint b) const - { - return vsub4(a, b); - } - - __host__ __device__ __forceinline__ VSub4() {} - __host__ __device__ __forceinline__ VSub4(const VSub4&) {} - }; - - struct VSub2 : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, uint b) const - { - return vsub2(a, b); - } - - __host__ __device__ __forceinline__ VSub2() {} - __host__ __device__ __forceinline__ VSub2(const VSub2&) {} - }; - - template struct SubMat : binary_function + template struct SubOp1 : binary_function { __device__ __forceinline__ D operator ()(T a, T b) const { return saturate_cast(a - b); } - - __host__ __device__ __forceinline__ SubMat() {} - __host__ __device__ __forceinline__ SubMat(const SubMat&) {} }; -} - -namespace cv { namespace cuda { namespace device -{ - template <> struct TransformFunctorTraits< arithm::VSub4 > : arithm::ArithmFuncTraits - { - }; - - template <> struct TransformFunctorTraits< arithm::VSub2 > : arithm::ArithmFuncTraits - { - }; - - template struct TransformFunctorTraits< arithm::SubMat > : arithm::ArithmFuncTraits - { - }; -}}} - -namespace arithm -{ - void subMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, VSub4(), WithOutMask(), stream); - } - - void subMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, VSub2(), WithOutMask(), stream); - } template - void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) + void subMat_v1(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, Stream& stream) { if (mask.data) - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, SubMat(), mask, stream); + gridTransformBinary(globPtr(src1), globPtr(src2), globPtr(dst), SubOp1(), globPtr(mask), stream); else - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, SubMat(), WithOutMask(), stream); + gridTransformBinary(globPtr(src1), globPtr(src2), globPtr(dst), SubOp1(), stream); } - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + struct SubOp2 : binary_function + { + __device__ __forceinline__ uint operator ()(uint a, uint b) const + { + return vsub2(a, b); + } + }; - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + void subMat_v2(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) + { + const int vcols = src1.cols >> 1; - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + GlobPtrSz src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols); + GlobPtrSz src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols); + GlobPtrSz dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + gridTransformBinary(src1_, src2_, dst_, SubOp2(), stream); + } - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + struct SubOp4 : binary_function + { + __device__ __forceinline__ uint operator ()(uint a, uint b) const + { + return vsub4(a, b); + } + }; - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + void subMat_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) + { + const int vcols = src1.cols >> 2; - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); + GlobPtrSz src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols); + GlobPtrSz src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols); + GlobPtrSz dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols); + + gridTransformBinary(src1_, src2_, dst_, SubOp4(), stream); + } } -#endif // CUDA_DISABLER +void subMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int) +{ + typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, Stream& stream); + static const func_t funcs[7][7] = + { + { + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1 + }, + { + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1 + }, + { + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1 + }, + { + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1, + subMat_v1 + }, + { + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + subMat_v1, + subMat_v1, + subMat_v1 + }, + { + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + subMat_v1, + subMat_v1 + }, + { + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + 0 /*subMat_v1*/, + subMat_v1 + } + }; + + const int sdepth = src1.depth(); + const int ddepth = dst.depth(); + + CV_DbgAssert( sdepth < 7 && ddepth < 7 ); + + GpuMat src1_ = src1.reshape(1); + GpuMat src2_ = src2.reshape(1); + GpuMat dst_ = dst.reshape(1); + + if (mask.empty() && (sdepth == CV_8U || sdepth == CV_16U) && ddepth == sdepth) + { + const intptr_t src1ptr = reinterpret_cast(src1_.data); + const intptr_t src2ptr = reinterpret_cast(src2_.data); + const intptr_t dstptr = reinterpret_cast(dst_.data); + + const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0; + + if (isAllAligned) + { + if (sdepth == CV_8U && (src1_.cols & 3) == 0) + { + subMat_v4(src1_, src2_, dst_, stream); + return; + } + else if (sdepth == CV_16U && (src1_.cols & 1) == 0) + { + subMat_v2(src1_, src2_, dst_, stream); + return; + } + } + } + + const func_t func = funcs[sdepth][ddepth]; + + if (!func) + CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); + + func(src1_, src2_, dst_, mask, stream); +} + +#endif diff --git a/modules/cudaarithm/src/cuda/sub_scalar.cu b/modules/cudaarithm/src/cuda/sub_scalar.cu index 5f4ef66a9c..35cea8cbe3 100644 --- a/modules/cudaarithm/src/cuda/sub_scalar.cu +++ b/modules/cudaarithm/src/cuda/sub_scalar.cu @@ -40,110 +40,164 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void subScalar(const GpuMat& src, cv::Scalar val, bool inv, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int); + +namespace { - template struct SubScalar : unary_function + template struct SubScalarOp : unary_function { - S val; - int scale; + ScalarType val; - __host__ SubScalar(S val_, int scale_) : val(val_), scale(scale_) {} - - __device__ __forceinline__ D operator ()(T a) const + __device__ __forceinline__ DstType operator ()(SrcType a) const { - return saturate_cast(scale * (a - val)); + return saturate_cast(saturate_cast(a) - val); } }; -} -namespace cv { namespace cuda { namespace device -{ - template struct TransformFunctorTraits< arithm::SubScalar > : arithm::ArithmFuncTraits + template struct SubScalarOpInv : unary_function + { + ScalarType val; + + __device__ __forceinline__ DstType operator ()(SrcType a) const + { + return saturate_cast(val - saturate_cast(a)); + } + }; + + template struct TransformPolicy : DefaultTransformPolicy { }; -}}} - -namespace arithm -{ - template - void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) + template <> struct TransformPolicy : DefaultTransformPolicy { - SubScalar op(static_cast(val), inv ? -1 : 1); + enum { + shift = 1 + }; + }; - if (mask.data) - device::transform((PtrStepSz) src1, (PtrStepSz) dst, op, mask, stream); + template + void subScalarImpl(const GpuMat& src, cv::Scalar value, bool inv, GpuMat& dst, const GpuMat& mask, Stream& stream) + { + typedef typename MakeVec::cn>::type ScalarType; + + cv::Scalar_ value_ = value; + + if (inv) + { + SubScalarOpInv op; + op.val = VecTraits::make(value_.val); + + if (mask.data) + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, globPtr(mask), stream); + else + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, stream); + } else - device::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + { + SubScalarOp op; + op.val = VecTraits::make(value_.val); + + if (mask.data) + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, globPtr(mask), stream); + else + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, stream); + } } - - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - //template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); } -#endif // CUDA_DISABLER +void subScalar(const GpuMat& src, cv::Scalar val, bool inv, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int) +{ + typedef void (*func_t)(const GpuMat& src, cv::Scalar val, bool inv, GpuMat& dst, const GpuMat& mask, Stream& stream); + static const func_t funcs[7][7][4] = + { + { + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl} + }, + { + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl} + }, + { + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl} + }, + { + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl} + }, + { + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl} + }, + { + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl} + }, + { + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/, 0 /*subScalarImpl*/}, + {subScalarImpl, subScalarImpl, subScalarImpl, subScalarImpl} + } + }; + + const int sdepth = src.depth(); + const int ddepth = dst.depth(); + const int cn = src.channels(); + + CV_DbgAssert( sdepth < 7 && ddepth < 7 && cn <= 4 ); + + const func_t func = funcs[sdepth][ddepth][cn - 1]; + + if (!func) + CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); + + func(src, val, inv, dst, mask, stream); +} + +#endif diff --git a/modules/cudaarithm/src/element_operations.cpp b/modules/cudaarithm/src/element_operations.cpp index b3711dcc1d..071c79368f 100644 --- a/modules/cudaarithm/src/element_operations.cpp +++ b/modules/cudaarithm/src/element_operations.cpp @@ -348,248 +348,9 @@ void cv::cuda::add(InputArray src1, InputArray src2, OutputArray dst, InputArray //////////////////////////////////////////////////////////////////////// // subtract -namespace arithm -{ - void subMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream); - void subMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream); +void subMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& _stream, int); - template - void subMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); -} - -static void subMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& _stream, int) -{ - typedef void (*func_t)(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - static const func_t funcs[7][7] = - { - { - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat - }, - { - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat - }, - { - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat - }, - { - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat, - arithm::subMat - }, - { - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - arithm::subMat, - arithm::subMat, - arithm::subMat - }, - { - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - arithm::subMat, - arithm::subMat - }, - { - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - 0 /*arithm::subMat*/, - arithm::subMat - } - }; - - const int sdepth = src1.depth(); - const int ddepth = dst.depth(); - const int cn = src1.channels(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - PtrStepSzb src1_(src1.rows, src1.cols * cn, src1.data, src1.step); - PtrStepSzb src2_(src1.rows, src1.cols * cn, src2.data, src2.step); - PtrStepSzb dst_(src1.rows, src1.cols * cn, dst.data, dst.step); - - if (mask.empty() && (sdepth == CV_8U || sdepth == CV_16U) && ddepth == sdepth) - { - const intptr_t src1ptr = reinterpret_cast(src1_.data); - const intptr_t src2ptr = reinterpret_cast(src2_.data); - const intptr_t dstptr = reinterpret_cast(dst_.data); - - const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0; - - if (isAllAligned) - { - if (sdepth == CV_8U && (src1_.cols & 3) == 0) - { - const int vcols = src1_.cols >> 2; - - arithm::subMat_v4(PtrStepSz(src1_.rows, vcols, (unsigned int*) src1_.data, src1_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) src2_.data, src2_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) dst_.data, dst_.step), - stream); - - return; - } - else if (sdepth == CV_16U && (src1_.cols & 1) == 0) - { - const int vcols = src1_.cols >> 1; - - arithm::subMat_v2(PtrStepSz(src1_.rows, vcols, (unsigned int*) src1_.data, src1_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) src2_.data, src2_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) dst_.data, dst_.step), - stream); - - return; - } - } - } - - const func_t func = funcs[sdepth][ddepth]; - - if (!func) - CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - - func(src1_, src2_, dst_, mask, stream); -} - -namespace arithm -{ - template - void subScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); -} - -static void subScalar(const GpuMat& src, Scalar val, bool inv, GpuMat& dst, const GpuMat& mask, double, Stream& _stream, int) -{ - typedef void (*func_t)(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - static const func_t funcs[7][7] = - { - { - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar - }, - { - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar - }, - { - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar - }, - { - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar - }, - { - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - arithm::subScalar, - arithm::subScalar, - arithm::subScalar - }, - { - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - arithm::subScalar, - arithm::subScalar - }, - { - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - 0 /*arithm::subScalar*/, - arithm::subScalar - } - }; - - typedef void (*npp_func_t)(const PtrStepSzb src, Scalar sc, PtrStepb dst, cudaStream_t stream); - static const npp_func_t npp_funcs[7][4] = - { - {NppArithmScalar::call, 0 , NppArithmScalar::call, NppArithmScalar::call}, - {0 , 0 , 0 , 0 }, - {NppArithmScalar::call, 0 , NppArithmScalar::call, NppArithmScalar::call}, - {NppArithmScalar::call, NppArithmScalar::call, NppArithmScalar::call, NppArithmScalar::call}, - {NppArithmScalar::call, NppArithmScalar::call, NppArithmScalar::call, 0 }, - {NppArithmScalar::call, NppArithmScalar::call, NppArithmScalar::call, NppArithmScalar::call}, - {0 , 0 , 0 , 0 } - }; - - const int sdepth = src.depth(); - const int ddepth = dst.depth(); - const int cn = src.channels(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - const npp_func_t npp_func = npp_funcs[sdepth][cn - 1]; - if (ddepth == sdepth && cn > 1 && npp_func != 0 && !inv) - { - npp_func(src, val, dst, stream); - return; - } - - CV_Assert( cn == 1 ); - - const func_t func = funcs[sdepth][ddepth]; - - if (!func) - CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - - func(src, val[0], inv, dst, mask, stream); -} +void subScalar(const GpuMat& src, Scalar val, bool inv, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int); void cv::cuda::subtract(InputArray src1, InputArray src2, OutputArray dst, InputArray mask, int dtype, Stream& stream) { From e7dba695b3941fd124438bce79800d508bddf9eb Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 29 Jul 2013 15:45:50 +0400 Subject: [PATCH 04/35] used new device layer for cv::gpu::multiply --- modules/cudaarithm/src/cuda/mul_mat.cu | 277 +++++++++--------- modules/cudaarithm/src/cuda/mul_scalar.cu | 208 +++++++------ modules/cudaarithm/src/element_operations.cpp | 220 +------------- 3 files changed, 274 insertions(+), 431 deletions(-) diff --git a/modules/cudaarithm/src/cuda/mul_mat.cu b/modules/cudaarithm/src/cuda/mul_mat.cu index dda596341a..f45e4e2094 100644 --- a/modules/cudaarithm/src/cuda/mul_mat.cu +++ b/modules/cudaarithm/src/cuda/mul_mat.cu @@ -40,172 +40,185 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void mulMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double scale, Stream& stream, int); +void mulMat_8uc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream); +void mulMat_16sc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream); + +namespace { - struct Mul_8uc4_32f : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, float b) const - { - uint res = 0; - - res |= (saturate_cast((0xffu & (a )) * b) ); - res |= (saturate_cast((0xffu & (a >> 8)) * b) << 8); - res |= (saturate_cast((0xffu & (a >> 16)) * b) << 16); - res |= (saturate_cast((0xffu & (a >> 24)) * b) << 24); - - return res; - } - - __host__ __device__ __forceinline__ Mul_8uc4_32f() {} - __host__ __device__ __forceinline__ Mul_8uc4_32f(const Mul_8uc4_32f&) {} - }; - - struct Mul_16sc4_32f : binary_function - { - __device__ __forceinline__ short4 operator ()(short4 a, float b) const - { - return make_short4(saturate_cast(a.x * b), saturate_cast(a.y * b), - saturate_cast(a.z * b), saturate_cast(a.w * b)); - } - - __host__ __device__ __forceinline__ Mul_16sc4_32f() {} - __host__ __device__ __forceinline__ Mul_16sc4_32f(const Mul_16sc4_32f&) {} - }; - - template struct Mul : binary_function + template struct MulOp : binary_function { __device__ __forceinline__ D operator ()(T a, T b) const { return saturate_cast(a * b); } - - __host__ __device__ __forceinline__ Mul() {} - __host__ __device__ __forceinline__ Mul(const Mul&) {} }; - template struct MulScale : binary_function + template struct MulScaleOp : binary_function { S scale; - __host__ explicit MulScale(S scale_) : scale(scale_) {} - __device__ __forceinline__ D operator ()(T a, T b) const { return saturate_cast(scale * a * b); } }; -} -namespace cv { namespace cuda { namespace device -{ - template <> struct TransformFunctorTraits : arithm::ArithmFuncTraits + template struct TransformPolicy : DefaultTransformPolicy { }; - - template struct TransformFunctorTraits< arithm::Mul > : arithm::ArithmFuncTraits + template <> struct TransformPolicy : DefaultTransformPolicy { + enum { + shift = 1 + }; }; - template struct TransformFunctorTraits< arithm::MulScale > : arithm::ArithmFuncTraits - { - }; -}}} - -namespace arithm -{ - void mulMat_8uc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, Mul_8uc4_32f(), WithOutMask(), stream); - } - - void mulMat_16sc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, Mul_16sc4_32f(), WithOutMask(), stream); - } - template - void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream) + void mulMatImpl(const GpuMat& src1, const GpuMat& src2, const GpuMat& dst, double scale, Stream& stream) { if (scale == 1) { - Mul op; - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + MulOp op; + gridTransformBinary_< TransformPolicy >(globPtr(src1), globPtr(src2), globPtr(dst), op, stream); } else { - MulScale op(static_cast(scale)); - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + MulScaleOp op; + op.scale = static_cast(scale); + gridTransformBinary_< TransformPolicy >(globPtr(src1), globPtr(src2), globPtr(dst), op, stream); } } - - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); } -#endif // CUDA_DISABLER +void mulMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double scale, Stream& stream, int) +{ + typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, const GpuMat& dst, double scale, Stream& stream); + static const func_t funcs[7][7] = + { + { + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl + }, + { + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl + }, + { + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl + }, + { + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl, + mulMatImpl + }, + { + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + mulMatImpl, + mulMatImpl, + mulMatImpl + }, + { + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + mulMatImpl, + mulMatImpl + }, + { + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + 0 /*mulMatImpl*/, + mulMatImpl + } + }; + + const int sdepth = src1.depth(); + const int ddepth = dst.depth(); + + CV_DbgAssert( sdepth < 7 && ddepth < 7 ); + + GpuMat src1_ = src1.reshape(1); + GpuMat src2_ = src2.reshape(1); + GpuMat dst_ = dst.reshape(1); + + const func_t func = funcs[sdepth][ddepth]; + + if (!func) + CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); + + func(src1_, src2_, dst_, scale, stream); +} + +namespace +{ + template + struct MulOpSpecial : binary_function + { + __device__ __forceinline__ T operator ()(const T& a, float b) const + { + typedef typename VecTraits::elem_type elem_type; + + T res; + + res.x = saturate_cast(a.x * b); + res.y = saturate_cast(a.y * b); + res.z = saturate_cast(a.z * b); + res.w = saturate_cast(a.w * b); + + return res; + } + }; +} + +void mulMat_8uc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) +{ + gridTransformBinary(globPtr(src1), globPtr(src2), globPtr(dst), MulOpSpecial(), stream); +} + +void mulMat_16sc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) +{ + gridTransformBinary(globPtr(src1), globPtr(src2), globPtr(dst), MulOpSpecial(), stream); +} + +#endif diff --git a/modules/cudaarithm/src/cuda/mul_scalar.cu b/modules/cudaarithm/src/cuda/mul_scalar.cu index 341cd9bc44..4700d30152 100644 --- a/modules/cudaarithm/src/cuda/mul_scalar.cu +++ b/modules/cudaarithm/src/cuda/mul_scalar.cu @@ -40,105 +40,143 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void mulScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMat& mask, double scale, Stream& stream, int); + +namespace { - template struct MulScalar : unary_function + template struct MulScalarOp : unary_function { - S val; + ScalarType val; - __host__ explicit MulScalar(S val_) : val(val_) {} - - __device__ __forceinline__ D operator ()(T a) const + __device__ __forceinline__ DstType operator ()(SrcType a) const { - return saturate_cast(a * val); + return saturate_cast(saturate_cast(a) * val); } }; -} -namespace cv { namespace cuda { namespace device -{ - template struct TransformFunctorTraits< arithm::MulScalar > : arithm::ArithmFuncTraits + template struct TransformPolicy : DefaultTransformPolicy { }; -}}} - -namespace arithm -{ - template - void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream) + template <> struct TransformPolicy : DefaultTransformPolicy { - MulScalar op(static_cast(val)); - device::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + enum { + shift = 1 + }; + }; + + template + void mulScalarImpl(const GpuMat& src, cv::Scalar value, GpuMat& dst, Stream& stream) + { + typedef typename MakeVec::cn>::type ScalarType; + + cv::Scalar_ value_ = value; + + MulScalarOp op; + op.val = VecTraits::make(value_.val); + + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, stream); } - - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - //template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - template void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); } -#endif // CUDA_DISABLER +void mulScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMat&, double scale, Stream& stream, int) +{ + typedef void (*func_t)(const GpuMat& src, cv::Scalar val, GpuMat& dst, Stream& stream); + static const func_t funcs[7][7][4] = + { + { + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl} + }, + { + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl} + }, + { + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl} + }, + { + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl} + }, + { + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl} + }, + { + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl} + }, + { + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/, 0 /*mulScalarImpl*/}, + {mulScalarImpl, mulScalarImpl, mulScalarImpl, mulScalarImpl} + } + }; + + const int sdepth = src.depth(); + const int ddepth = dst.depth(); + const int cn = src.channels(); + + CV_DbgAssert( sdepth < 7 && ddepth < 7 && cn <= 4 ); + + val[0] *= scale; + val[1] *= scale; + val[2] *= scale; + val[3] *= scale; + + const func_t func = funcs[sdepth][ddepth][cn - 1]; + + if (!func) + CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); + + func(src, val, dst, stream); +} + +#endif diff --git a/modules/cudaarithm/src/element_operations.cpp b/modules/cudaarithm/src/element_operations.cpp index 071c79368f..d170166a7a 100644 --- a/modules/cudaarithm/src/element_operations.cpp +++ b/modules/cudaarithm/src/element_operations.cpp @@ -360,219 +360,11 @@ void cv::cuda::subtract(InputArray src1, InputArray src2, OutputArray dst, Input //////////////////////////////////////////////////////////////////////// // multiply -namespace arithm -{ - void mulMat_8uc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream); +void mulMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double scale, Stream& stream, int); +void mulMat_8uc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream); +void mulMat_16sc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream); - void mulMat_16sc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream); - - template - void mulMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); -} - -static void mulMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double scale, Stream& _stream, int) -{ - typedef void (*func_t)(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - static const func_t funcs[7][7] = - { - { - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat - }, - { - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat - }, - { - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat - }, - { - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat - }, - { - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - arithm::mulMat, - arithm::mulMat, - arithm::mulMat - }, - { - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - arithm::mulMat, - arithm::mulMat - }, - { - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - 0 /*arithm::mulMat*/, - arithm::mulMat - } - }; - - const int sdepth = src1.depth(); - const int ddepth = dst.depth(); - const int cn = src1.channels(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - PtrStepSzb src1_(src1.rows, src1.cols * cn, src1.data, src1.step); - PtrStepSzb src2_(src1.rows, src1.cols * cn, src2.data, src2.step); - PtrStepSzb dst_(src1.rows, src1.cols * cn, dst.data, dst.step); - - const func_t func = funcs[sdepth][ddepth]; - - if (!func) - CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - - func(src1_, src2_, dst_, scale, stream); -} - -namespace arithm -{ - template - void mulScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); -} - -static void mulScalar(const GpuMat& src, Scalar val, bool, GpuMat& dst, const GpuMat&, double scale, Stream& _stream, int) -{ - typedef void (*func_t)(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[7][7] = - { - { - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar - }, - { - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar - }, - { - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar - }, - { - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar - }, - { - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - arithm::mulScalar, - arithm::mulScalar, - arithm::mulScalar - }, - { - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - arithm::mulScalar, - arithm::mulScalar - }, - { - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - 0 /*arithm::mulScalar*/, - arithm::mulScalar - } - }; - - typedef void (*npp_func_t)(const PtrStepSzb src, Scalar sc, PtrStepb dst, cudaStream_t stream); - static const npp_func_t npp_funcs[7][4] = - { - {NppArithmScalar::call, 0, NppArithmScalar::call, NppArithmScalar::call}, - {0 , 0, 0 , 0 }, - {NppArithmScalar::call, 0, NppArithmScalar::call, NppArithmScalar::call}, - {NppArithmScalar::call, 0, NppArithmScalar::call, NppArithmScalar::call}, - {NppArithmScalar::call, 0, NppArithmScalar::call, 0 }, - {NppArithmScalar::call, 0, NppArithmScalar::call, NppArithmScalar::call}, - {0 , 0, 0 , 0 } - }; - - const int sdepth = src.depth(); - const int ddepth = dst.depth(); - const int cn = src.channels(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - val[0] *= scale; - val[1] *= scale; - val[2] *= scale; - val[3] *= scale; - - const npp_func_t npp_func = npp_funcs[sdepth][cn - 1]; - if (ddepth == sdepth && cn > 1 && npp_func != 0) - { - npp_func(src, val, dst, stream); - return; - } - - CV_Assert( cn == 1 ); - - const func_t func = funcs[sdepth][ddepth]; - - if (!func) - CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - - func(src, val[0], dst, stream); -} +void mulScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMat& mask, double scale, Stream& stream, int); void cv::cuda::multiply(InputArray _src1, InputArray _src2, OutputArray _dst, double scale, int dtype, Stream& stream) { @@ -586,7 +378,7 @@ void cv::cuda::multiply(InputArray _src1, InputArray _src2, OutputArray _dst, do _dst.create(src1.size(), src1.type()); GpuMat dst = _dst.getGpuMat(); - arithm::mulMat_8uc4_32f(src1, src2, dst, StreamAccessor::getStream(stream)); + mulMat_8uc4_32f(src1, src2, dst, stream); } else if (_src1.type() == CV_16SC4 && _src2.type() == CV_32FC1) { @@ -598,7 +390,7 @@ void cv::cuda::multiply(InputArray _src1, InputArray _src2, OutputArray _dst, do _dst.create(src1.size(), src1.type()); GpuMat dst = _dst.getGpuMat(); - arithm::mulMat_16sc4_32f(src1, src2, dst, StreamAccessor::getStream(stream)); + mulMat_16sc4_32f(src1, src2, dst, stream); } else { From 574ff471466606b504402e95f95ea219437e34b2 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 29 Jul 2013 16:18:27 +0400 Subject: [PATCH 05/35] used new device layer for cv::gpu::divide --- modules/cudaarithm/src/cuda/div_mat.cu | 294 +++++++++--------- modules/cudaarithm/src/cuda/div_scalar.cu | 286 +++++++++++------ modules/cudaarithm/src/element_operations.cpp | 230 +------------- 3 files changed, 350 insertions(+), 460 deletions(-) diff --git a/modules/cudaarithm/src/cuda/div_mat.cu b/modules/cudaarithm/src/cuda/div_mat.cu index 4a62a44ae3..e139cb4ae9 100644 --- a/modules/cudaarithm/src/cuda/div_mat.cu +++ b/modules/cudaarithm/src/cuda/div_mat.cu @@ -40,191 +40,203 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void divMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double scale, Stream& stream, int); +void divMat_8uc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream); +void divMat_16sc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream); + +namespace { - struct Div_8uc4_32f : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, float b) const - { - uint res = 0; - - if (b != 0) - { - b = 1.0f / b; - res |= (saturate_cast((0xffu & (a )) * b) ); - res |= (saturate_cast((0xffu & (a >> 8)) * b) << 8); - res |= (saturate_cast((0xffu & (a >> 16)) * b) << 16); - res |= (saturate_cast((0xffu & (a >> 24)) * b) << 24); - } - - return res; - } - }; - - struct Div_16sc4_32f : binary_function - { - __device__ __forceinline__ short4 operator ()(short4 a, float b) const - { - return b != 0 ? make_short4(saturate_cast(a.x / b), saturate_cast(a.y / b), - saturate_cast(a.z / b), saturate_cast(a.w / b)) - : make_short4(0,0,0,0); - } - }; - - template struct Div : binary_function + template struct DivOp : binary_function { __device__ __forceinline__ D operator ()(T a, T b) const { return b != 0 ? saturate_cast(a / b) : 0; } - - __host__ __device__ __forceinline__ Div() {} - __host__ __device__ __forceinline__ Div(const Div&) {} }; - template struct Div : binary_function + template struct DivOp : binary_function { __device__ __forceinline__ float operator ()(T a, T b) const { - return b != 0 ? static_cast(a) / b : 0; + return b != 0 ? static_cast(a) / b : 0.0f; } - - __host__ __device__ __forceinline__ Div() {} - __host__ __device__ __forceinline__ Div(const Div&) {} }; - template struct Div : binary_function + template struct DivOp : binary_function { __device__ __forceinline__ double operator ()(T a, T b) const { - return b != 0 ? static_cast(a) / b : 0; + return b != 0 ? static_cast(a) / b : 0.0; } - - __host__ __device__ __forceinline__ Div() {} - __host__ __device__ __forceinline__ Div(const Div&) {} }; - template struct DivScale : binary_function + template struct DivScaleOp : binary_function { S scale; - __host__ explicit DivScale(S scale_) : scale(scale_) {} - __device__ __forceinline__ D operator ()(T a, T b) const { return b != 0 ? saturate_cast(scale * a / b) : 0; } }; -} -namespace cv { namespace cuda { namespace device -{ - template <> struct TransformFunctorTraits : arithm::ArithmFuncTraits + template struct TransformPolicy : DefaultTransformPolicy { }; - - template struct TransformFunctorTraits< arithm::Div > : arithm::ArithmFuncTraits + template <> struct TransformPolicy : DefaultTransformPolicy { + enum { + shift = 1 + }; }; - template struct TransformFunctorTraits< arithm::DivScale > : arithm::ArithmFuncTraits - { - }; -}}} - -namespace arithm -{ - void divMat_8uc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, Div_8uc4_32f(), WithOutMask(), stream); - } - - void divMat_16sc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, Div_16sc4_32f(), WithOutMask(), stream); - } - template - void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream) + void divMatImpl(const GpuMat& src1, const GpuMat& src2, const GpuMat& dst, double scale, Stream& stream) { if (scale == 1) { - Div op; - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + DivOp op; + gridTransformBinary_< TransformPolicy >(globPtr(src1), globPtr(src2), globPtr(dst), op, stream); } else { - DivScale op(static_cast(scale)); - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + DivScaleOp op; + op.scale = static_cast(scale); + gridTransformBinary_< TransformPolicy >(globPtr(src1), globPtr(src2), globPtr(dst), op, stream); } } - - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - //template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - template void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); } -#endif // CUDA_DISABLER +void divMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double scale, Stream& stream, int) +{ + typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, const GpuMat& dst, double scale, Stream& stream); + static const func_t funcs[7][7] = + { + { + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl + }, + { + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl + }, + { + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl + }, + { + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl, + divMatImpl + }, + { + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + divMatImpl, + divMatImpl, + divMatImpl + }, + { + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + divMatImpl, + divMatImpl + }, + { + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + 0 /*divMatImpl*/, + divMatImpl + } + }; + + const int sdepth = src1.depth(); + const int ddepth = dst.depth(); + + CV_DbgAssert( sdepth < 7 && ddepth < 7 ); + + GpuMat src1_ = src1.reshape(1); + GpuMat src2_ = src2.reshape(1); + GpuMat dst_ = dst.reshape(1); + + const func_t func = funcs[sdepth][ddepth]; + + if (!func) + CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); + + func(src1_, src2_, dst_, scale, stream); +} + +namespace +{ + template + struct DivOpSpecial : binary_function + { + __device__ __forceinline__ T operator ()(const T& a, float b) const + { + typedef typename VecTraits::elem_type elem_type; + + T res = VecTraits::all(0); + + if (b != 0) + { + b = 1.0f / b; + res.x = saturate_cast(a.x * b); + res.y = saturate_cast(a.y * b); + res.z = saturate_cast(a.z * b); + res.w = saturate_cast(a.w * b); + } + + return res; + } + }; +} + +void divMat_8uc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) +{ + gridTransformBinary(globPtr(src1), globPtr(src2), globPtr(dst), DivOpSpecial(), stream); +} + +void divMat_16sc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) +{ + gridTransformBinary(globPtr(src1), globPtr(src2), globPtr(dst), DivOpSpecial(), stream); +} + +#endif diff --git a/modules/cudaarithm/src/cuda/div_scalar.cu b/modules/cudaarithm/src/cuda/div_scalar.cu index 9ee1727a75..186e1766e9 100644 --- a/modules/cudaarithm/src/cuda/div_scalar.cu +++ b/modules/cudaarithm/src/cuda/div_scalar.cu @@ -40,129 +40,225 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void divScalar(const GpuMat& src, cv::Scalar val, bool inv, GpuMat& dst, const GpuMat& mask, double scale, Stream& stream, int); + +namespace { - template struct DivScalar : unary_function + template struct SafeDiv; + template <> struct SafeDiv<1> { - S val; - - __host__ explicit DivScalar(S val_) : val(val_) {} - - __device__ __forceinline__ D operator ()(T a) const + template + __device__ __forceinline__ static T op(T a, T b) { - return saturate_cast(a / val); + return b != 0 ? a / b : 0; + } + }; + template <> struct SafeDiv<2> + { + template + __device__ __forceinline__ static T op(const T& a, const T& b) + { + T res; + + res.x = b.x != 0 ? a.x / b.x : 0; + res.y = b.y != 0 ? a.y / b.y : 0; + + return res; + } + }; + template <> struct SafeDiv<3> + { + template + __device__ __forceinline__ static T op(const T& a, const T& b) + { + T res; + + res.x = b.x != 0 ? a.x / b.x : 0; + res.y = b.y != 0 ? a.y / b.y : 0; + res.z = b.z != 0 ? a.z / b.z : 0; + + return res; + } + }; + template <> struct SafeDiv<4> + { + template + __device__ __forceinline__ static T op(const T& a, const T& b) + { + T res; + + res.x = b.x != 0 ? a.x / b.x : 0; + res.y = b.y != 0 ? a.y / b.y : 0; + res.z = b.z != 0 ? a.z / b.z : 0; + res.w = b.w != 0 ? a.w / b.w : 0; + + return res; } }; - template struct DivScalarInv : unary_function + template struct DivScalarOp : unary_function { - S val; + ScalarType val; - explicit DivScalarInv(S val_) : val(val_) {} - - __device__ __forceinline__ D operator ()(T a) const + __device__ __forceinline__ DstType operator ()(SrcType a) const { - return a != 0 ? saturate_cast(val / a) : 0; + return saturate_cast(SafeDiv::cn>::op(saturate_cast(a), val)); } }; -} -namespace cv { namespace cuda { namespace device -{ - template struct TransformFunctorTraits< arithm::DivScalar > : arithm::ArithmFuncTraits + template struct DivScalarOpInv : unary_function { + ScalarType val; + + __device__ __forceinline__ DstType operator ()(SrcType a) const + { + return saturate_cast(SafeDiv::cn>::op(val, saturate_cast(a))); + } }; - template struct TransformFunctorTraits< arithm::DivScalarInv > : arithm::ArithmFuncTraits + template struct TransformPolicy : DefaultTransformPolicy { }; -}}} - -namespace arithm -{ - template - void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream) + template <> struct TransformPolicy : DefaultTransformPolicy { + enum { + shift = 1 + }; + }; + + template + void divScalarImpl(const GpuMat& src, cv::Scalar value, bool inv, GpuMat& dst, Stream& stream) + { + typedef typename MakeVec::cn>::type ScalarType; + + cv::Scalar_ value_ = value; + if (inv) { - DivScalarInv op(static_cast(val)); - device::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + DivScalarOpInv op; + op.val = VecTraits::make(value_.val); + + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, stream); } else { - DivScalar op(static_cast(val)); - device::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + DivScalarOp op; + op.val = VecTraits::make(value_.val); + + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, stream); } } - - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - //template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - template void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); } -#endif // CUDA_DISABLER +void divScalar(const GpuMat& src, cv::Scalar val, bool inv, GpuMat& dst, const GpuMat&, double scale, Stream& stream, int) +{ + typedef void (*func_t)(const GpuMat& src, cv::Scalar val, bool inv, GpuMat& dst, Stream& stream); + static const func_t funcs[7][7][4] = + { + { + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl} + }, + { + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl} + }, + { + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl} + }, + { + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl} + }, + { + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl} + }, + { + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl} + }, + { + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/, 0 /*divScalarImpl*/}, + {divScalarImpl, divScalarImpl, divScalarImpl, divScalarImpl} + } + }; + + const int sdepth = src.depth(); + const int ddepth = dst.depth(); + const int cn = src.channels(); + + CV_DbgAssert( sdepth < 7 && ddepth < 7 && cn <= 4 ); + + if (inv) + { + val[0] *= scale; + val[1] *= scale; + val[2] *= scale; + val[3] *= scale; + } + else + { + val[0] /= scale; + val[1] /= scale; + val[2] /= scale; + val[3] /= scale; + } + + const func_t func = funcs[sdepth][ddepth][cn - 1]; + + if (!func) + CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); + + func(src, val, inv, dst, stream); +} + +#endif diff --git a/modules/cudaarithm/src/element_operations.cpp b/modules/cudaarithm/src/element_operations.cpp index d170166a7a..b846520e86 100644 --- a/modules/cudaarithm/src/element_operations.cpp +++ b/modules/cudaarithm/src/element_operations.cpp @@ -401,229 +401,11 @@ void cv::cuda::multiply(InputArray _src1, InputArray _src2, OutputArray _dst, do //////////////////////////////////////////////////////////////////////// // divide -namespace arithm -{ - void divMat_8uc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream); +void divMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double scale, Stream& stream, int); +void divMat_8uc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream); +void divMat_16sc4_32f(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream); - void divMat_16sc4_32f(PtrStepSz src1, PtrStepSzf src2, PtrStepSz dst, cudaStream_t stream); - - template - void divMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); -} - -static void divMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double scale, Stream& _stream, int) -{ - typedef void (*func_t)(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, double scale, cudaStream_t stream); - static const func_t funcs[7][7] = - { - { - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat - }, - { - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat - }, - { - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat - }, - { - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat, - arithm::divMat - }, - { - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - arithm::divMat, - arithm::divMat, - arithm::divMat - }, - { - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - arithm::divMat, - arithm::divMat - }, - { - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - 0 /*arithm::divMat*/, - arithm::divMat - } - }; - - const int sdepth = src1.depth(); - const int ddepth = dst.depth(); - const int cn = src1.channels(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - PtrStepSzb src1_(src1.rows, src1.cols * cn, src1.data, src1.step); - PtrStepSzb src2_(src1.rows, src1.cols * cn, src2.data, src2.step); - PtrStepSzb dst_(src1.rows, src1.cols * cn, dst.data, dst.step); - - const func_t func = funcs[sdepth][ddepth]; - - if (!func) - CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - - func(src1_, src2_, dst_, scale, stream); -} - -namespace arithm -{ - template - void divScalar(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); -} - -static void divScalar(const GpuMat& src, Scalar val, bool inv, GpuMat& dst, const GpuMat&, double scale, Stream& _stream, int) -{ - typedef void (*func_t)(PtrStepSzb src1, double val, bool inv, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[7][7] = - { - { - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar - }, - { - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar - }, - { - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar - }, - { - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar - }, - { - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - arithm::divScalar, - arithm::divScalar, - arithm::divScalar - }, - { - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - arithm::divScalar, - arithm::divScalar - }, - { - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - 0 /*arithm::divScalar*/, - arithm::divScalar - } - }; - - typedef void (*npp_func_t)(const PtrStepSzb src, Scalar sc, PtrStepb dst, cudaStream_t stream); - static const npp_func_t npp_funcs[7][4] = - { - {NppArithmScalar::call, 0, NppArithmScalar::call, NppArithmScalar::call}, - {0 , 0, 0 , 0 }, - {NppArithmScalar::call, 0, NppArithmScalar::call, NppArithmScalar::call}, - {NppArithmScalar::call, 0, NppArithmScalar::call, NppArithmScalar::call}, - {NppArithmScalar::call, 0, NppArithmScalar::call, 0 }, - {NppArithmScalar::call, 0, NppArithmScalar::call, NppArithmScalar::call}, - {0 , 0, 0 , 0 } - }; - - const int sdepth = src.depth(); - const int ddepth = dst.depth(); - const int cn = src.channels(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - if (inv) - { - val[0] *= scale; - val[1] *= scale; - val[2] *= scale; - val[3] *= scale; - } - else - { - val[0] /= scale; - val[1] /= scale; - val[2] /= scale; - val[3] /= scale; - } - - const npp_func_t npp_func = npp_funcs[sdepth][cn - 1]; - if (ddepth == sdepth && cn > 1 && npp_func != 0 && !inv) - { - npp_func(src, val, dst, stream); - return; - } - - CV_Assert( cn == 1 ); - - const func_t func = funcs[sdepth][ddepth]; - - if (!func) - CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - - func(src, val[0], inv, dst, stream); -} +void divScalar(const GpuMat& src, cv::Scalar val, bool inv, GpuMat& dst, const GpuMat& mask, double scale, Stream& stream, int); void cv::cuda::divide(InputArray _src1, InputArray _src2, OutputArray _dst, double scale, int dtype, Stream& stream) { @@ -637,7 +419,7 @@ void cv::cuda::divide(InputArray _src1, InputArray _src2, OutputArray _dst, doub _dst.create(src1.size(), src1.type()); GpuMat dst = _dst.getGpuMat(); - arithm::divMat_8uc4_32f(src1, src2, dst, StreamAccessor::getStream(stream)); + divMat_8uc4_32f(src1, src2, dst, stream); } else if (_src1.type() == CV_16SC4 && _src2.type() == CV_32FC1) { @@ -649,7 +431,7 @@ void cv::cuda::divide(InputArray _src1, InputArray _src2, OutputArray _dst, doub _dst.create(src1.size(), src1.type()); GpuMat dst = _dst.getGpuMat(); - arithm::divMat_16sc4_32f(src1, src2, dst, StreamAccessor::getStream(stream)); + divMat_16sc4_32f(src1, src2, dst, stream); } else { From 7628e57fc62a46b788790c698bd7e06aba92d515 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 29 Jul 2013 16:43:37 +0400 Subject: [PATCH 06/35] used new device layer for cv::gpu::absdiff --- modules/cudaarithm/src/cuda/absdiff_mat.cu | 179 +++++++++++------- modules/cudaarithm/src/cuda/absdiff_scalar.cu | 80 ++++---- modules/cudaarithm/src/element_operations.cpp | 100 +--------- 3 files changed, 158 insertions(+), 201 deletions(-) diff --git a/modules/cudaarithm/src/cuda/absdiff_mat.cu b/modules/cudaarithm/src/cuda/absdiff_mat.cu index ffdcbfa380..0f93b00eba 100644 --- a/modules/cudaarithm/src/cuda/absdiff_mat.cu +++ b/modules/cudaarithm/src/cuda/absdiff_mat.cu @@ -40,43 +40,22 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void absDiffMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int); + +namespace { - struct VAbsDiff4 : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, uint b) const - { - return vabsdiff4(a, b); - } - - __host__ __device__ __forceinline__ VAbsDiff4() {} - __host__ __device__ __forceinline__ VAbsDiff4(const VAbsDiff4&) {} - }; - - struct VAbsDiff2 : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, uint b) const - { - return vabsdiff2(a, b); - } - - __host__ __device__ __forceinline__ VAbsDiff2() {} - __host__ __device__ __forceinline__ VAbsDiff2(const VAbsDiff2&) {} - }; - __device__ __forceinline__ int _abs(int a) { return ::abs(a); @@ -90,58 +69,120 @@ namespace arithm return ::fabs(a); } - template struct AbsDiffMat : binary_function + template struct AbsDiffOp1 : binary_function { __device__ __forceinline__ T operator ()(T a, T b) const { return saturate_cast(_abs(a - b)); } - - __host__ __device__ __forceinline__ AbsDiffMat() {} - __host__ __device__ __forceinline__ AbsDiffMat(const AbsDiffMat&) {} - }; -} - -namespace cv { namespace cuda { namespace device -{ - template <> struct TransformFunctorTraits< arithm::VAbsDiff4 > : arithm::ArithmFuncTraits - { }; - template <> struct TransformFunctorTraits< arithm::VAbsDiff2 > : arithm::ArithmFuncTraits + template struct TransformPolicy : DefaultTransformPolicy { }; - - template struct TransformFunctorTraits< arithm::AbsDiffMat > : arithm::ArithmFuncTraits + template <> struct TransformPolicy : DefaultTransformPolicy { + enum { + shift = 1 + }; }; -}}} - -namespace arithm -{ - void absDiffMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, VAbsDiff4(), WithOutMask(), stream); - } - - void absDiffMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, VAbsDiff2(), WithOutMask(), stream); - } template - void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream) + void absDiffMat_v1(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) { - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, AbsDiffMat(), WithOutMask(), stream); + gridTransformBinary_< TransformPolicy >(globPtr(src1), globPtr(src2), globPtr(dst), AbsDiffOp1(), stream); } - template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); + struct AbsDiffOp2 : binary_function + { + __device__ __forceinline__ uint operator ()(uint a, uint b) const + { + return vabsdiff2(a, b); + } + }; + + void absDiffMat_v2(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) + { + const int vcols = src1.cols >> 1; + + GlobPtrSz src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols); + GlobPtrSz src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols); + GlobPtrSz dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols); + + gridTransformBinary(src1_, src2_, dst_, AbsDiffOp2(), stream); + } + + struct AbsDiffOp4 : binary_function + { + __device__ __forceinline__ uint operator ()(uint a, uint b) const + { + return vabsdiff4(a, b); + } + }; + + void absDiffMat_v4(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream) + { + const int vcols = src1.cols >> 2; + + GlobPtrSz src1_ = globPtr((uint*) src1.data, src1.step, src1.rows, vcols); + GlobPtrSz src2_ = globPtr((uint*) src2.data, src2.step, src1.rows, vcols); + GlobPtrSz dst_ = globPtr((uint*) dst.data, dst.step, src1.rows, vcols); + + gridTransformBinary(src1_, src2_, dst_, AbsDiffOp4(), stream); + } } -#endif // CUDA_DISABLER +void absDiffMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int) +{ + typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream); + static const func_t funcs[] = + { + absDiffMat_v1, + absDiffMat_v1, + absDiffMat_v1, + absDiffMat_v1, + absDiffMat_v1, + absDiffMat_v1, + absDiffMat_v1 + }; + + const int depth = src1.depth(); + + CV_DbgAssert( depth < 7 ); + + GpuMat src1_ = src1.reshape(1); + GpuMat src2_ = src2.reshape(1); + GpuMat dst_ = dst.reshape(1); + + if (depth == CV_8U || depth == CV_16U) + { + const intptr_t src1ptr = reinterpret_cast(src1_.data); + const intptr_t src2ptr = reinterpret_cast(src2_.data); + const intptr_t dstptr = reinterpret_cast(dst_.data); + + const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0; + + if (isAllAligned) + { + if (depth == CV_8U && (src1_.cols & 3) == 0) + { + absDiffMat_v4(src1_, src2_, dst_, stream); + return; + } + else if (depth == CV_16U && (src1_.cols & 1) == 0) + { + absDiffMat_v2(src1_, src2_, dst_, stream); + return; + } + } + } + + const func_t func = funcs[depth]; + + if (!func) + CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); + + func(src1_, src2_, dst_, stream); +} + +#endif diff --git a/modules/cudaarithm/src/cuda/absdiff_scalar.cu b/modules/cudaarithm/src/cuda/absdiff_scalar.cu index ad30bffda6..f6cebdac77 100644 --- a/modules/cudaarithm/src/cuda/absdiff_scalar.cu +++ b/modules/cudaarithm/src/cuda/absdiff_scalar.cu @@ -40,59 +40,71 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void absDiffScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMat&, double, Stream& stream, int); + +namespace { - template struct AbsDiffScalar : unary_function + template struct AbsDiffScalarOp : unary_function { S val; - __host__ explicit AbsDiffScalar(S val_) : val(val_) {} - __device__ __forceinline__ T operator ()(T a) const { abs_func f; return saturate_cast(f(a - val)); } }; -} -namespace cv { namespace cuda { namespace device -{ - template struct TransformFunctorTraits< arithm::AbsDiffScalar > : arithm::ArithmFuncTraits + template struct TransformPolicy : DefaultTransformPolicy { }; -}}} - -namespace arithm -{ - template - void absDiffScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream) + template <> struct TransformPolicy : DefaultTransformPolicy { - AbsDiffScalar op(static_cast(val)); + enum { + shift = 1 + }; + }; - device::transform((PtrStepSz) src1, (PtrStepSz) dst, op, WithOutMask(), stream); + template + void absDiffScalarImpl(const GpuMat& src, double value, GpuMat& dst, Stream& stream) + { + AbsDiffScalarOp op; + op.val = static_cast(value); + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, stream); } - - template void absDiffScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); - template void absDiffScalar(PtrStepSzb src1, double src2, PtrStepSzb dst, cudaStream_t stream); } -#endif // CUDA_DISABLER +void absDiffScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMat&, double, Stream& stream, int) +{ + typedef void (*func_t)(const GpuMat& src, double val, GpuMat& dst, Stream& stream); + static const func_t funcs[] = + { + absDiffScalarImpl, + absDiffScalarImpl, + absDiffScalarImpl, + absDiffScalarImpl, + absDiffScalarImpl, + absDiffScalarImpl, + absDiffScalarImpl + }; + + const int depth = src.depth(); + + CV_DbgAssert( depth < 7 ); + + funcs[depth](src, val[0], dst, stream); +} + +#endif diff --git a/modules/cudaarithm/src/element_operations.cpp b/modules/cudaarithm/src/element_operations.cpp index b846520e86..e6700619f1 100644 --- a/modules/cudaarithm/src/element_operations.cpp +++ b/modules/cudaarithm/src/element_operations.cpp @@ -442,105 +442,9 @@ void cv::cuda::divide(InputArray _src1, InputArray _src2, OutputArray _dst, doub ////////////////////////////////////////////////////////////////////////////// // absdiff -namespace arithm -{ - void absDiffMat_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream); - void absDiffMat_v2(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream); +void absDiffMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int); - template - void absDiffMat(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); -} - -static void absDiffMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& _stream, int) -{ - typedef void (*func_t)(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[] = - { - arithm::absDiffMat, - arithm::absDiffMat, - arithm::absDiffMat, - arithm::absDiffMat, - arithm::absDiffMat, - arithm::absDiffMat, - arithm::absDiffMat - }; - - const int depth = src1.depth(); - const int cn = src1.channels(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - PtrStepSzb src1_(src1.rows, src1.cols * cn, src1.data, src1.step); - PtrStepSzb src2_(src1.rows, src1.cols * cn, src2.data, src2.step); - PtrStepSzb dst_(src1.rows, src1.cols * cn, dst.data, dst.step); - - if (depth == CV_8U || depth == CV_16U) - { - const intptr_t src1ptr = reinterpret_cast(src1_.data); - const intptr_t src2ptr = reinterpret_cast(src2_.data); - const intptr_t dstptr = reinterpret_cast(dst_.data); - - const bool isAllAligned = (src1ptr & 31) == 0 && (src2ptr & 31) == 0 && (dstptr & 31) == 0; - - if (isAllAligned) - { - if (depth == CV_8U && (src1_.cols & 3) == 0) - { - const int vcols = src1_.cols >> 2; - - arithm::absDiffMat_v4(PtrStepSz(src1_.rows, vcols, (unsigned int*) src1_.data, src1_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) src2_.data, src2_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) dst_.data, dst_.step), - stream); - - return; - } - else if (depth == CV_16U && (src1_.cols & 1) == 0) - { - const int vcols = src1_.cols >> 1; - - arithm::absDiffMat_v2(PtrStepSz(src1_.rows, vcols, (unsigned int*) src1_.data, src1_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) src2_.data, src2_.step), - PtrStepSz(src1_.rows, vcols, (unsigned int*) dst_.data, dst_.step), - stream); - - return; - } - } - } - - const func_t func = funcs[depth]; - - if (!func) - CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - - func(src1_, src2_, dst_, stream); -} - -namespace arithm -{ - template - void absDiffScalar(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); -} - -static void absDiffScalar(const GpuMat& src, Scalar val, bool, GpuMat& dst, const GpuMat&, double, Stream& stream, int) -{ - typedef void (*func_t)(PtrStepSzb src1, double val, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[] = - { - arithm::absDiffScalar, - arithm::absDiffScalar, - arithm::absDiffScalar, - arithm::absDiffScalar, - arithm::absDiffScalar, - arithm::absDiffScalar, - arithm::absDiffScalar - }; - - const int depth = src.depth(); - - funcs[depth](src, val[0], dst, StreamAccessor::getStream(stream)); -} +void absDiffScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMat&, double, Stream& stream, int); void cv::cuda::absdiff(InputArray src1, InputArray src2, OutputArray dst, Stream& stream) { From 766d950ff32223525bda4df330d55e567099feb0 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 26 Aug 2013 10:11:29 +0400 Subject: [PATCH 07/35] used new device layer in math per element operations --- modules/cudaarithm/src/cuda/math.cu | 359 ++++++++++-------- modules/cudaarithm/src/element_operations.cpp | 264 ------------- 2 files changed, 208 insertions(+), 415 deletions(-) diff --git a/modules/cudaarithm/src/cuda/math.cu b/modules/cudaarithm/src/cuda/math.cu index ecd9a8cf90..39f822081d 100644 --- a/modules/cudaarithm/src/cuda/math.cu +++ b/modules/cudaarithm/src/cuda/math.cu @@ -40,196 +40,248 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" -#include "opencv2/core/cuda/limits.hpp" -#include "opencv2/core/cuda/type_traits.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -////////////////////////////////////////////////////////////////////////// -// absMat +#include "opencv2/cudaarithm.hpp" +#include "opencv2/cudev.hpp" -namespace cv { namespace cuda { namespace device +using namespace cv::cudev; + +namespace { - template struct TransformFunctorTraits< abs_func > : arithm::ArithmFuncTraits + template struct TransformPolicy : DefaultTransformPolicy { }; -}}} - -namespace arithm -{ - template - void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) + template <> struct TransformPolicy : DefaultTransformPolicy { - device::transform((PtrStepSz) src, (PtrStepSz) dst, abs_func(), WithOutMask(), stream); - } - - template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); + enum { + shift = 1 + }; + }; } -////////////////////////////////////////////////////////////////////////// -// sqrMat +////////////////////////////////////////////////////////////////////////////// +/// abs -namespace arithm +namespace { - template struct Sqr : unary_function + template + void absMat(const GpuMat& src, const GpuMat& dst, Stream& stream) + { + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), abs_func(), stream); + } +} + +void cv::cuda::abs(InputArray _src, OutputArray _dst, Stream& stream) +{ + typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream); + static const func_t funcs[] = + { + absMat, + absMat, + absMat, + absMat, + absMat, + absMat, + absMat + }; + + GpuMat src = _src.getGpuMat(); + + const int depth = src.depth(); + + CV_DbgAssert( depth <= CV_64F ); + + _dst.create(src.size(), src.type()); + GpuMat dst = _dst.getGpuMat(); + + funcs[depth](src.reshape(1), dst.reshape(1), stream); +} + +////////////////////////////////////////////////////////////////////////////// +/// sqr + +namespace +{ + template struct SqrOp : unary_function { __device__ __forceinline__ T operator ()(T x) const { return saturate_cast(x * x); } - - __host__ __device__ __forceinline__ Sqr() {} - __host__ __device__ __forceinline__ Sqr(const Sqr&) {} }; + + template + void sqrMat(const GpuMat& src, const GpuMat& dst, Stream& stream) + { + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), SqrOp(), stream); + } } -namespace cv { namespace cuda { namespace device +void cv::cuda::sqr(InputArray _src, OutputArray _dst, Stream& stream) { - template struct TransformFunctorTraits< arithm::Sqr > : arithm::ArithmFuncTraits + typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream); + static const func_t funcs[] = { + sqrMat, + sqrMat, + sqrMat, + sqrMat, + sqrMat, + sqrMat, + sqrMat }; -}}} -namespace arithm + GpuMat src = _src.getGpuMat(); + + const int depth = src.depth(); + + CV_DbgAssert( depth <= CV_64F ); + + _dst.create(src.size(), src.type()); + GpuMat dst = _dst.getGpuMat(); + + funcs[depth](src.reshape(1), dst.reshape(1), stream); +} + +////////////////////////////////////////////////////////////////////////////// +/// sqrt + +namespace { template - void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) + void sqrtMat(const GpuMat& src, const GpuMat& dst, Stream& stream) { - device::transform((PtrStepSz) src, (PtrStepSz) dst, Sqr(), WithOutMask(), stream); + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), sqrt_func(), stream); } - - template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); } -////////////////////////////////////////////////////////////////////////// -// sqrtMat - -namespace cv { namespace cuda { namespace device +void cv::cuda::sqrt(InputArray _src, OutputArray _dst, Stream& stream) { - template struct TransformFunctorTraits< sqrt_func > : arithm::ArithmFuncTraits + typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream); + static const func_t funcs[] = { + sqrtMat, + sqrtMat, + sqrtMat, + sqrtMat, + sqrtMat, + sqrtMat, + sqrtMat }; -}}} -namespace arithm -{ - template - void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) - { - device::transform((PtrStepSz) src, (PtrStepSz) dst, sqrt_func(), WithOutMask(), stream); - } + GpuMat src = _src.getGpuMat(); - template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); + const int depth = src.depth(); + + CV_DbgAssert( depth <= CV_64F ); + + _dst.create(src.size(), src.type()); + GpuMat dst = _dst.getGpuMat(); + + funcs[depth](src.reshape(1), dst.reshape(1), stream); } -////////////////////////////////////////////////////////////////////////// -// logMat +//////////////////////////////////////////////////////////////////////// +/// exp -namespace cv { namespace cuda { namespace device +namespace { - template struct TransformFunctorTraits< log_func > : arithm::ArithmFuncTraits - { - }; -}}} - -namespace arithm -{ - template - void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) - { - device::transform((PtrStepSz) src, (PtrStepSz) dst, log_func(), WithOutMask(), stream); - } - - template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); -} - -////////////////////////////////////////////////////////////////////////// -// expMat - -namespace arithm -{ - template struct Exp : unary_function + template struct ExpOp : unary_function { __device__ __forceinline__ T operator ()(T x) const { exp_func f; return saturate_cast(f(x)); } - - __host__ __device__ __forceinline__ Exp() {} - __host__ __device__ __forceinline__ Exp(const Exp&) {} }; + + template + void expMat(const GpuMat& src, const GpuMat& dst, Stream& stream) + { + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), ExpOp(), stream); + } } -namespace cv { namespace cuda { namespace device +void cv::cuda::exp(InputArray _src, OutputArray _dst, Stream& stream) { - template struct TransformFunctorTraits< arithm::Exp > : arithm::ArithmFuncTraits + typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream); + static const func_t funcs[] = { + expMat, + expMat, + expMat, + expMat, + expMat, + expMat, + expMat }; -}}} -namespace arithm + GpuMat src = _src.getGpuMat(); + + const int depth = src.depth(); + + CV_DbgAssert( depth <= CV_64F ); + + _dst.create(src.size(), src.type()); + GpuMat dst = _dst.getGpuMat(); + + funcs[depth](src.reshape(1), dst.reshape(1), stream); +} + +//////////////////////////////////////////////////////////////////////// +// log + +namespace { template - void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream) + void logMat(const GpuMat& src, const GpuMat& dst, Stream& stream) { - device::transform((PtrStepSz) src, (PtrStepSz) dst, Exp(), WithOutMask(), stream); + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), log_func(), stream); } - - template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - template void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); } -////////////////////////////////////////////////////////////////////////// +void cv::cuda::log(InputArray _src, OutputArray _dst, Stream& stream) +{ + typedef void (*func_t)(const GpuMat& src, const GpuMat& dst, Stream& stream); + static const func_t funcs[] = + { + logMat, + logMat, + logMat, + logMat, + logMat, + logMat, + logMat + }; + + GpuMat src = _src.getGpuMat(); + + const int depth = src.depth(); + + CV_DbgAssert( depth <= CV_64F ); + + _dst.create(src.size(), src.type()); + GpuMat dst = _dst.getGpuMat(); + + funcs[depth](src.reshape(1), dst.reshape(1), stream); +} + +//////////////////////////////////////////////////////////////////////// // pow -namespace arithm +namespace { template::is_signed> struct PowOp : unary_function { float power; - __host__ explicit PowOp(double power_) : power(static_cast(power_)) {} - __device__ __forceinline__ T operator()(T e) const { return saturate_cast(__powf((float)e, power)); @@ -239,8 +291,6 @@ namespace arithm { float power; - __host__ explicit PowOp(double power_) : power(static_cast(power_)) {} - __device__ __forceinline__ T operator()(T e) const { T res = saturate_cast(__powf((float)e, power)); @@ -255,8 +305,6 @@ namespace arithm { float power; - __host__ explicit PowOp(double power_) : power(static_cast(power_)) {} - __device__ __forceinline__ float operator()(float e) const { return __powf(::fabs(e), power); @@ -266,37 +314,46 @@ namespace arithm { double power; - __host__ explicit PowOp(double power_) : power(power_) {} - __device__ __forceinline__ double operator()(double e) const { return ::pow(::fabs(e), power); } }; -} -namespace cv { namespace cuda { namespace device -{ - template struct TransformFunctorTraits< arithm::PowOp > : arithm::ArithmFuncTraits - { - }; -}}} - -namespace arithm -{ template - void pow(PtrStepSzb src, double power, PtrStepSzb dst, cudaStream_t stream) + void powMat(const GpuMat& src, double power, const GpuMat& dst, Stream& stream) { - device::transform((PtrStepSz) src, (PtrStepSz) dst, PowOp(power), WithOutMask(), stream); - } + PowOp op; + op.power = static_cast::type>(power); - template void pow(PtrStepSzb src, double power, PtrStepSzb dst, cudaStream_t stream); - template void pow(PtrStepSzb src, double power, PtrStepSzb dst, cudaStream_t stream); - template void pow(PtrStepSzb src, double power, PtrStepSzb dst, cudaStream_t stream); - template void pow(PtrStepSzb src, double power, PtrStepSzb dst, cudaStream_t stream); - template void pow(PtrStepSzb src, double power, PtrStepSzb dst, cudaStream_t stream); - template void pow(PtrStepSzb src, double power, PtrStepSzb dst, cudaStream_t stream); - template void pow(PtrStepSzb src, double power, PtrStepSzb dst, cudaStream_t stream); + gridTransformUnary_< TransformPolicy >(globPtr(src), globPtr(dst), op, stream); + } } -#endif // CUDA_DISABLER +void cv::cuda::pow(InputArray _src, double power, OutputArray _dst, Stream& stream) +{ + typedef void (*func_t)(const GpuMat& src, double power, const GpuMat& dst, Stream& stream); + static const func_t funcs[] = + { + powMat, + powMat, + powMat, + powMat, + powMat, + powMat, + powMat + }; + + GpuMat src = _src.getGpuMat(); + + const int depth = src.depth(); + + CV_DbgAssert(depth <= CV_64F); + + _dst.create(src.size(), src.type()); + GpuMat dst = _dst.getGpuMat(); + + funcs[depth](src.reshape(1), power, dst.reshape(1), stream); +} + +#endif diff --git a/modules/cudaarithm/src/element_operations.cpp b/modules/cudaarithm/src/element_operations.cpp index e6700619f1..c3bbced9e5 100644 --- a/modules/cudaarithm/src/element_operations.cpp +++ b/modules/cudaarithm/src/element_operations.cpp @@ -451,270 +451,6 @@ void cv::cuda::absdiff(InputArray src1, InputArray src2, OutputArray dst, Stream arithm_op(src1, src2, dst, noArray(), 1.0, -1, stream, absDiffMat, absDiffScalar); } -////////////////////////////////////////////////////////////////////////////// -// abs - -namespace arithm -{ - template - void absMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); -} - -void cv::cuda::abs(InputArray _src, OutputArray _dst, Stream& stream) -{ - using namespace arithm; - - typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[] = - { - absMat, - absMat, - absMat, - absMat, - absMat, - absMat, - absMat - }; - - GpuMat src = _src.getGpuMat(); - - const int depth = src.depth(); - - CV_Assert( depth <= CV_64F ); - CV_Assert( src.channels() == 1 ); - - if (depth == CV_64F) - { - if (!deviceSupports(NATIVE_DOUBLE)) - CV_Error(cv::Error::StsUnsupportedFormat, "The device doesn't support double"); - } - - _dst.create(src.size(), src.type()); - GpuMat dst = _dst.getGpuMat(); - - funcs[depth](src, dst, StreamAccessor::getStream(stream)); -} - -////////////////////////////////////////////////////////////////////////////// -// sqr - -namespace arithm -{ - template - void sqrMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); -} - -void cv::cuda::sqr(InputArray _src, OutputArray _dst, Stream& stream) -{ - using namespace arithm; - - typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[] = - { - sqrMat, - sqrMat, - sqrMat, - sqrMat, - sqrMat, - sqrMat, - sqrMat - }; - - GpuMat src = _src.getGpuMat(); - - const int depth = src.depth(); - - CV_Assert( depth <= CV_64F ); - CV_Assert( src.channels() == 1 ); - - if (depth == CV_64F) - { - if (!deviceSupports(NATIVE_DOUBLE)) - CV_Error(cv::Error::StsUnsupportedFormat, "The device doesn't support double"); - } - - _dst.create(src.size(), src.type()); - GpuMat dst = _dst.getGpuMat(); - - funcs[depth](src, dst, StreamAccessor::getStream(stream)); -} - -////////////////////////////////////////////////////////////////////////////// -// sqrt - -namespace arithm -{ - template - void sqrtMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); -} - -void cv::cuda::sqrt(InputArray _src, OutputArray _dst, Stream& stream) -{ - using namespace arithm; - - typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[] = - { - sqrtMat, - sqrtMat, - sqrtMat, - sqrtMat, - sqrtMat, - sqrtMat, - sqrtMat - }; - - GpuMat src = _src.getGpuMat(); - - const int depth = src.depth(); - - CV_Assert( depth <= CV_64F ); - CV_Assert( src.channels() == 1 ); - - if (depth == CV_64F) - { - if (!deviceSupports(NATIVE_DOUBLE)) - CV_Error(cv::Error::StsUnsupportedFormat, "The device doesn't support double"); - } - - _dst.create(src.size(), src.type()); - GpuMat dst = _dst.getGpuMat(); - - funcs[depth](src, dst, StreamAccessor::getStream(stream)); -} - -//////////////////////////////////////////////////////////////////////// -// exp - -namespace arithm -{ - template - void expMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); -} - -void cv::cuda::exp(InputArray _src, OutputArray _dst, Stream& stream) -{ - using namespace arithm; - - typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[] = - { - expMat, - expMat, - expMat, - expMat, - expMat, - expMat, - expMat - }; - - GpuMat src = _src.getGpuMat(); - - const int depth = src.depth(); - - CV_Assert( depth <= CV_64F ); - CV_Assert( src.channels() == 1 ); - - if (depth == CV_64F) - { - if (!deviceSupports(NATIVE_DOUBLE)) - CV_Error(cv::Error::StsUnsupportedFormat, "The device doesn't support double"); - } - - _dst.create(src.size(), src.type()); - GpuMat dst = _dst.getGpuMat(); - - funcs[depth](src, dst, StreamAccessor::getStream(stream)); -} - -//////////////////////////////////////////////////////////////////////// -// log - -namespace arithm -{ - template - void logMat(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); -} - -void cv::cuda::log(InputArray _src, OutputArray _dst, Stream& stream) -{ - using namespace arithm; - - typedef void (*func_t)(PtrStepSzb src, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[] = - { - logMat, - logMat, - logMat, - logMat, - logMat, - logMat, - logMat - }; - - GpuMat src = _src.getGpuMat(); - - const int depth = src.depth(); - - CV_Assert( depth <= CV_64F ); - CV_Assert( src.channels() == 1 ); - - if (depth == CV_64F) - { - if (!deviceSupports(NATIVE_DOUBLE)) - CV_Error(cv::Error::StsUnsupportedFormat, "The device doesn't support double"); - } - - _dst.create(src.size(), src.type()); - GpuMat dst = _dst.getGpuMat(); - - funcs[depth](src, dst, StreamAccessor::getStream(stream)); -} - -//////////////////////////////////////////////////////////////////////// -// pow - -namespace arithm -{ - template void pow(PtrStepSzb src, double power, PtrStepSzb dst, cudaStream_t stream); -} - -void cv::cuda::pow(InputArray _src, double power, OutputArray _dst, Stream& stream) -{ - typedef void (*func_t)(PtrStepSzb src, double power, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[] = - { - arithm::pow, - arithm::pow, - arithm::pow, - arithm::pow, - arithm::pow, - arithm::pow, - arithm::pow - }; - - GpuMat src = _src.getGpuMat(); - - const int depth = src.depth(); - const int cn = src.channels(); - - CV_Assert(depth <= CV_64F); - - if (depth == CV_64F) - { - if (!deviceSupports(NATIVE_DOUBLE)) - CV_Error(cv::Error::StsUnsupportedFormat, "The device doesn't support double"); - } - - _dst.create(src.size(), src.type()); - GpuMat dst = _dst.getGpuMat(); - - PtrStepSzb src_(src.rows, src.cols * cn, src.data, src.step); - PtrStepSzb dst_(src.rows, src.cols * cn, dst.data, dst.step); - - funcs[depth](src_, power, dst_, StreamAccessor::getStream(stream)); -} - ////////////////////////////////////////////////////////////////////////////// // compare From 43e811118bb44e12e217ef4c144d791db6c25022 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 29 Jul 2013 17:13:43 +0400 Subject: [PATCH 08/35] fixed assertion --- modules/cudaarithm/src/cuda/absdiff_mat.cu | 2 +- modules/cudaarithm/src/cuda/absdiff_scalar.cu | 2 +- modules/cudaarithm/src/cuda/add_mat.cu | 2 +- modules/cudaarithm/src/cuda/add_scalar.cu | 2 +- modules/cudaarithm/src/cuda/div_mat.cu | 2 +- modules/cudaarithm/src/cuda/div_scalar.cu | 2 +- modules/cudaarithm/src/cuda/mul_mat.cu | 2 +- modules/cudaarithm/src/cuda/mul_scalar.cu | 2 +- modules/cudaarithm/src/cuda/sub_mat.cu | 2 +- modules/cudaarithm/src/cuda/sub_scalar.cu | 2 +- 10 files changed, 10 insertions(+), 10 deletions(-) diff --git a/modules/cudaarithm/src/cuda/absdiff_mat.cu b/modules/cudaarithm/src/cuda/absdiff_mat.cu index 0f93b00eba..ec04f12284 100644 --- a/modules/cudaarithm/src/cuda/absdiff_mat.cu +++ b/modules/cudaarithm/src/cuda/absdiff_mat.cu @@ -148,7 +148,7 @@ void absDiffMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMa const int depth = src1.depth(); - CV_DbgAssert( depth < 7 ); + CV_DbgAssert( depth <= CV_64F ); GpuMat src1_ = src1.reshape(1); GpuMat src2_ = src2.reshape(1); diff --git a/modules/cudaarithm/src/cuda/absdiff_scalar.cu b/modules/cudaarithm/src/cuda/absdiff_scalar.cu index f6cebdac77..3ffd0661b7 100644 --- a/modules/cudaarithm/src/cuda/absdiff_scalar.cu +++ b/modules/cudaarithm/src/cuda/absdiff_scalar.cu @@ -102,7 +102,7 @@ void absDiffScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const G const int depth = src.depth(); - CV_DbgAssert( depth < 7 ); + CV_DbgAssert( depth <= CV_64F ); funcs[depth](src, val[0], dst, stream); } diff --git a/modules/cudaarithm/src/cuda/add_mat.cu b/modules/cudaarithm/src/cuda/add_mat.cu index 6e7a7925fd..4166cc104e 100644 --- a/modules/cudaarithm/src/cuda/add_mat.cu +++ b/modules/cudaarithm/src/cuda/add_mat.cu @@ -185,7 +185,7 @@ void addMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& m const int sdepth = src1.depth(); const int ddepth = dst.depth(); - CV_DbgAssert( sdepth < 7 && ddepth < 7 ); + CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F ); GpuMat src1_ = src1.reshape(1); GpuMat src2_ = src2.reshape(1); diff --git a/modules/cudaarithm/src/cuda/add_scalar.cu b/modules/cudaarithm/src/cuda/add_scalar.cu index e0788e9bdd..92838a2a57 100644 --- a/modules/cudaarithm/src/cuda/add_scalar.cu +++ b/modules/cudaarithm/src/cuda/add_scalar.cu @@ -167,7 +167,7 @@ void addScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMa const int ddepth = dst.depth(); const int cn = src.channels(); - CV_DbgAssert( sdepth < 7 && ddepth < 7 && cn <= 4 ); + CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F && cn <= 4 ); const func_t func = funcs[sdepth][ddepth][cn - 1]; diff --git a/modules/cudaarithm/src/cuda/div_mat.cu b/modules/cudaarithm/src/cuda/div_mat.cu index e139cb4ae9..2a2fb9bf51 100644 --- a/modules/cudaarithm/src/cuda/div_mat.cu +++ b/modules/cudaarithm/src/cuda/div_mat.cu @@ -190,7 +190,7 @@ void divMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, const int sdepth = src1.depth(); const int ddepth = dst.depth(); - CV_DbgAssert( sdepth < 7 && ddepth < 7 ); + CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F ); GpuMat src1_ = src1.reshape(1); GpuMat src2_ = src2.reshape(1); diff --git a/modules/cudaarithm/src/cuda/div_scalar.cu b/modules/cudaarithm/src/cuda/div_scalar.cu index 186e1766e9..d3d2ce03f5 100644 --- a/modules/cudaarithm/src/cuda/div_scalar.cu +++ b/modules/cudaarithm/src/cuda/div_scalar.cu @@ -236,7 +236,7 @@ void divScalar(const GpuMat& src, cv::Scalar val, bool inv, GpuMat& dst, const G const int ddepth = dst.depth(); const int cn = src.channels(); - CV_DbgAssert( sdepth < 7 && ddepth < 7 && cn <= 4 ); + CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F && cn <= 4 ); if (inv) { diff --git a/modules/cudaarithm/src/cuda/mul_mat.cu b/modules/cudaarithm/src/cuda/mul_mat.cu index f45e4e2094..6ea7065573 100644 --- a/modules/cudaarithm/src/cuda/mul_mat.cu +++ b/modules/cudaarithm/src/cuda/mul_mat.cu @@ -176,7 +176,7 @@ void mulMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, const int sdepth = src1.depth(); const int ddepth = dst.depth(); - CV_DbgAssert( sdepth < 7 && ddepth < 7 ); + CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F ); GpuMat src1_ = src1.reshape(1); GpuMat src2_ = src2.reshape(1); diff --git a/modules/cudaarithm/src/cuda/mul_scalar.cu b/modules/cudaarithm/src/cuda/mul_scalar.cu index 4700d30152..f27ef26ddd 100644 --- a/modules/cudaarithm/src/cuda/mul_scalar.cu +++ b/modules/cudaarithm/src/cuda/mul_scalar.cu @@ -164,7 +164,7 @@ void mulScalar(const GpuMat& src, cv::Scalar val, bool, GpuMat& dst, const GpuMa const int ddepth = dst.depth(); const int cn = src.channels(); - CV_DbgAssert( sdepth < 7 && ddepth < 7 && cn <= 4 ); + CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F && cn <= 4 ); val[0] *= scale; val[1] *= scale; diff --git a/modules/cudaarithm/src/cuda/sub_mat.cu b/modules/cudaarithm/src/cuda/sub_mat.cu index ec8d229229..6468692aee 100644 --- a/modules/cudaarithm/src/cuda/sub_mat.cu +++ b/modules/cudaarithm/src/cuda/sub_mat.cu @@ -185,7 +185,7 @@ void subMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& m const int sdepth = src1.depth(); const int ddepth = dst.depth(); - CV_DbgAssert( sdepth < 7 && ddepth < 7 ); + CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F ); GpuMat src1_ = src1.reshape(1); GpuMat src2_ = src2.reshape(1); diff --git a/modules/cudaarithm/src/cuda/sub_scalar.cu b/modules/cudaarithm/src/cuda/sub_scalar.cu index 35cea8cbe3..c4eeec0148 100644 --- a/modules/cudaarithm/src/cuda/sub_scalar.cu +++ b/modules/cudaarithm/src/cuda/sub_scalar.cu @@ -190,7 +190,7 @@ void subScalar(const GpuMat& src, cv::Scalar val, bool inv, GpuMat& dst, const G const int ddepth = dst.depth(); const int cn = src.channels(); - CV_DbgAssert( sdepth < 7 && ddepth < 7 && cn <= 4 ); + CV_DbgAssert( sdepth <= CV_64F && ddepth <= CV_64F && cn <= 4 ); const func_t func = funcs[sdepth][ddepth][cn - 1]; From ef9917ecf141604063070cca5195bad4f58f015c Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 29 Jul 2013 17:53:03 +0400 Subject: [PATCH 09/35] used new device layer for cv::gpu::compare --- modules/cudaarithm/src/cuda/cmp_mat.cu | 275 +++++++-------- modules/cudaarithm/src/cuda/cmp_scalar.cu | 317 +++++++----------- modules/cudaarithm/src/element_operations.cpp | 142 +------- .../include/opencv2/cudev/util/vec_traits.hpp | 2 +- 4 files changed, 276 insertions(+), 460 deletions(-) diff --git a/modules/cudaarithm/src/cuda/cmp_mat.cu b/modules/cudaarithm/src/cuda/cmp_mat.cu index cdbb9639b4..3693fc2b78 100644 --- a/modules/cudaarithm/src/cuda/cmp_mat.cu +++ b/modules/cudaarithm/src/cuda/cmp_mat.cu @@ -40,64 +40,23 @@ // //M*/ -#if !defined CUDA_DISABLER +#include "opencv2/opencv_modules.hpp" -#include "opencv2/core/cuda/common.hpp" -#include "opencv2/core/cuda/functional.hpp" -#include "opencv2/core/cuda/transform.hpp" -#include "opencv2/core/cuda/saturate_cast.hpp" -#include "opencv2/core/cuda/simd_functions.hpp" +#ifndef HAVE_OPENCV_CUDEV -#include "arithm_func_traits.hpp" +#error "opencv_cudev is required" -using namespace cv::cuda; -using namespace cv::cuda::device; +#else -namespace arithm +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void cmpMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat&, double, Stream& stream, int cmpop); + +namespace { - struct VCmpEq4 : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, uint b) const - { - return vcmpeq4(a, b); - } - - __host__ __device__ __forceinline__ VCmpEq4() {} - __host__ __device__ __forceinline__ VCmpEq4(const VCmpEq4&) {} - }; - struct VCmpNe4 : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, uint b) const - { - return vcmpne4(a, b); - } - - __host__ __device__ __forceinline__ VCmpNe4() {} - __host__ __device__ __forceinline__ VCmpNe4(const VCmpNe4&) {} - }; - struct VCmpLt4 : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, uint b) const - { - return vcmplt4(a, b); - } - - __host__ __device__ __forceinline__ VCmpLt4() {} - __host__ __device__ __forceinline__ VCmpLt4(const VCmpLt4&) {} - }; - struct VCmpLe4 : binary_function - { - __device__ __forceinline__ uint operator ()(uint a, uint b) const - { - return vcmple4(a, b); - } - - __host__ __device__ __forceinline__ VCmpLe4() {} - __host__ __device__ __forceinline__ VCmpLe4(const VCmpLe4&) {} - }; - - template - struct Cmp : binary_function + template struct CmpOp : binary_function { __device__ __forceinline__ uchar operator()(T a, T b) const { @@ -105,102 +64,156 @@ namespace arithm return -op(a, b); } }; -} -namespace cv { namespace cuda { namespace device -{ - template <> struct TransformFunctorTraits< arithm::VCmpEq4 > : arithm::ArithmFuncTraits + template struct TransformPolicy : DefaultTransformPolicy { }; - template <> struct TransformFunctorTraits< arithm::VCmpNe4 > : arithm::ArithmFuncTraits + template <> struct TransformPolicy : DefaultTransformPolicy { + enum { + shift = 1 + }; }; - template <> struct TransformFunctorTraits< arithm::VCmpLt4 > : arithm::ArithmFuncTraits - { - }; - template <> struct TransformFunctorTraits< arithm::VCmpLe4 > : arithm::ArithmFuncTraits - { - }; - - template struct TransformFunctorTraits< arithm::Cmp > : arithm::ArithmFuncTraits - { - }; -}}} - -namespace arithm -{ - void cmpMatEq_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, VCmpEq4(), WithOutMask(), stream); - } - void cmpMatNe_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, VCmpNe4(), WithOutMask(), stream); - } - void cmpMatLt_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, VCmpLt4(), WithOutMask(), stream); - } - void cmpMatLe_v4(PtrStepSz src1, PtrStepSz src2, PtrStepSz dst, cudaStream_t stream) - { - device::transform(src1, src2, dst, VCmpLe4(), WithOutMask(), stream); - } template