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/cudaarithm/CMakeLists.txt b/modules/cudaarithm/CMakeLists.txt index 67a7ff922f..f44c283c09 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) ocv_module_include_directories() ocv_glob_module_sources() diff --git a/modules/cudaarithm/perf/perf_arithm.cpp b/modules/cudaarithm/perf/perf_arithm.cpp index d0f3e66176..9004155016 100644 --- a/modules/cudaarithm/perf/perf_arithm.cpp +++ b/modules/cudaarithm/perf/perf_arithm.cpp @@ -248,60 +248,3 @@ PERF_TEST_P(Sz_KernelSz_Ccorr, Convolve, CPU_SANITY_CHECK(dst); } } - -////////////////////////////////////////////////////////////////////// -// Integral - -PERF_TEST_P(Sz, Integral, - CUDA_TYPICAL_MAT_SIZES) -{ - const cv::Size size = GetParam(); - - cv::Mat src(size, CV_8UC1); - declare.in(src, WARMUP_RNG); - - if (PERF_RUN_CUDA()) - { - const cv::cuda::GpuMat d_src(src); - cv::cuda::GpuMat dst; - cv::cuda::GpuMat d_buf; - - TEST_CYCLE() cv::cuda::integral(d_src, dst, d_buf); - - CUDA_SANITY_CHECK(dst); - } - else - { - cv::Mat dst; - - TEST_CYCLE() cv::integral(src, dst); - - CPU_SANITY_CHECK(dst); - } -} - -////////////////////////////////////////////////////////////////////// -// IntegralSqr - -PERF_TEST_P(Sz, IntegralSqr, - CUDA_TYPICAL_MAT_SIZES) -{ - const cv::Size size = GetParam(); - - cv::Mat src(size, CV_8UC1); - declare.in(src, WARMUP_RNG); - - if (PERF_RUN_CUDA()) - { - const cv::cuda::GpuMat d_src(src); - cv::cuda::GpuMat dst, buf; - - TEST_CYCLE() cv::cuda::sqrIntegral(d_src, dst, buf); - - CUDA_SANITY_CHECK(dst); - } - else - { - FAIL_NO_CPU(); - } -} diff --git a/modules/cudaarithm/perf/perf_reductions.cpp b/modules/cudaarithm/perf/perf_reductions.cpp index fe72795560..470df48a3f 100644 --- a/modules/cudaarithm/perf/perf_reductions.cpp +++ b/modules/cudaarithm/perf/perf_reductions.cpp @@ -373,7 +373,7 @@ PERF_TEST_P(Sz_Depth_Cn_Code_Dim, Reduce, const cv::cuda::GpuMat d_src(src); cv::cuda::GpuMat dst; - TEST_CYCLE() cv::cuda::reduce(d_src, dst, dim, reduceOp); + TEST_CYCLE() cv::cuda::reduce(d_src, dst, dim, reduceOp, CV_32F); CUDA_SANITY_CHECK(dst); } @@ -381,7 +381,7 @@ PERF_TEST_P(Sz_Depth_Cn_Code_Dim, Reduce, { cv::Mat dst; - TEST_CYCLE() cv::reduce(src, dst, dim, reduceOp); + TEST_CYCLE() cv::reduce(src, dst, dim, reduceOp, CV_32F); CPU_SANITY_CHECK(dst); } @@ -465,3 +465,60 @@ PERF_TEST_P(Sz, MeanStdDev, SANITY_CHECK(cpu_stddev); } } + +////////////////////////////////////////////////////////////////////// +// Integral + +PERF_TEST_P(Sz, Integral, + CUDA_TYPICAL_MAT_SIZES) +{ + const cv::Size size = GetParam(); + + cv::Mat src(size, CV_8UC1); + declare.in(src, WARMUP_RNG); + + if (PERF_RUN_CUDA()) + { + const cv::cuda::GpuMat d_src(src); + cv::cuda::GpuMat dst; + cv::cuda::GpuMat d_buf; + + TEST_CYCLE() cv::cuda::integral(d_src, dst, d_buf); + + CUDA_SANITY_CHECK(dst); + } + else + { + cv::Mat dst; + + TEST_CYCLE() cv::integral(src, dst); + + CPU_SANITY_CHECK(dst); + } +} + +////////////////////////////////////////////////////////////////////// +// IntegralSqr + +PERF_TEST_P(Sz, IntegralSqr, + CUDA_TYPICAL_MAT_SIZES) +{ + const cv::Size size = GetParam(); + + cv::Mat src(size, CV_8UC1); + declare.in(src, WARMUP_RNG); + + if (PERF_RUN_CUDA()) + { + const cv::cuda::GpuMat d_src(src); + cv::cuda::GpuMat dst, buf; + + TEST_CYCLE() cv::cuda::sqrIntegral(d_src, dst, buf); + + CUDA_SANITY_CHECK(dst); + } + else + { + FAIL_NO_CPU(); + } +} diff --git a/modules/cudaarithm/src/arithm.cpp b/modules/cudaarithm/src/arithm.cpp index 6f7417ae35..4fe38e935e 100644 --- a/modules/cudaarithm/src/arithm.cpp +++ b/modules/cudaarithm/src/arithm.cpp @@ -292,95 +292,6 @@ void cv::cuda::gemm(InputArray _src1, InputArray _src2, double alpha, InputArray #endif } -////////////////////////////////////////////////////////////////////////////// -// mulSpectrums - -#ifdef HAVE_CUFFT - -namespace cv { namespace cuda { namespace device -{ - void mulSpectrums(const PtrStep a, const PtrStep b, PtrStepSz c, cudaStream_t stream); - - void mulSpectrums_CONJ(const PtrStep a, const PtrStep b, PtrStepSz c, cudaStream_t stream); -}}} - -#endif - -void cv::cuda::mulSpectrums(InputArray _src1, InputArray _src2, OutputArray _dst, int flags, bool conjB, Stream& stream) -{ -#ifndef HAVE_CUFFT - (void) _src1; - (void) _src2; - (void) _dst; - (void) flags; - (void) conjB; - (void) stream; - throw_no_cuda(); -#else - (void) flags; - - typedef void (*Caller)(const PtrStep, const PtrStep, PtrStepSz, cudaStream_t stream); - static Caller callers[] = { device::mulSpectrums, device::mulSpectrums_CONJ }; - - GpuMat src1 = _src1.getGpuMat(); - GpuMat src2 = _src2.getGpuMat(); - - CV_Assert( src1.type() == src2.type() && src1.type() == CV_32FC2 ); - CV_Assert( src1.size() == src2.size() ); - - _dst.create(src1.size(), CV_32FC2); - GpuMat dst = _dst.getGpuMat(); - - Caller caller = callers[(int)conjB]; - caller(src1, src2, dst, StreamAccessor::getStream(stream)); -#endif -} - -////////////////////////////////////////////////////////////////////////////// -// mulAndScaleSpectrums - -#ifdef HAVE_CUFFT - -namespace cv { namespace cuda { namespace device -{ - void mulAndScaleSpectrums(const PtrStep a, const PtrStep b, float scale, PtrStepSz c, cudaStream_t stream); - - void mulAndScaleSpectrums_CONJ(const PtrStep a, const PtrStep b, float scale, PtrStepSz c, cudaStream_t stream); -}}} - -#endif - -void cv::cuda::mulAndScaleSpectrums(InputArray _src1, InputArray _src2, OutputArray _dst, int flags, float scale, bool conjB, Stream& stream) -{ -#ifndef HAVE_CUFFT - (void) _src1; - (void) _src2; - (void) _dst; - (void) flags; - (void) scale; - (void) conjB; - (void) stream; - throw_no_cuda(); -#else - (void)flags; - - typedef void (*Caller)(const PtrStep, const PtrStep, float scale, PtrStepSz, cudaStream_t stream); - static Caller callers[] = { device::mulAndScaleSpectrums, device::mulAndScaleSpectrums_CONJ }; - - GpuMat src1 = _src1.getGpuMat(); - GpuMat src2 = _src2.getGpuMat(); - - CV_Assert( src1.type() == src2.type() && src1.type() == CV_32FC2); - CV_Assert( src1.size() == src2.size() ); - - _dst.create(src1.size(), CV_32FC2); - GpuMat dst = _dst.getGpuMat(); - - Caller caller = callers[(int)conjB]; - caller(src1, src2, scale, dst, StreamAccessor::getStream(stream)); -#endif -} - ////////////////////////////////////////////////////////////////////////////// // dft diff --git a/modules/cudaarithm/src/core.cpp b/modules/cudaarithm/src/core.cpp index e6c14db09f..eb71d6a4ec 100644 --- a/modules/cudaarithm/src/core.cpp +++ b/modules/cudaarithm/src/core.cpp @@ -63,163 +63,6 @@ void cv::cuda::copyMakeBorder(InputArray, OutputArray, int, int, int, int, int, #else /* !defined (HAVE_CUDA) */ -//////////////////////////////////////////////////////////////////////// -// merge/split - -namespace cv { namespace cuda { namespace device -{ - namespace split_merge - { - void merge(const PtrStepSzb* src, PtrStepSzb& dst, int total_channels, size_t elem_size, const cudaStream_t& stream); - void split(const PtrStepSzb& src, PtrStepSzb* dst, int num_channels, size_t elem_size1, const cudaStream_t& stream); - } -}}} - -namespace -{ - void merge_caller(const GpuMat* src, size_t n, OutputArray _dst, Stream& stream) - { - CV_Assert( src != 0 ); - CV_Assert( n > 0 && n <= 4 ); - - const int depth = src[0].depth(); - const Size size = src[0].size(); - - for (size_t i = 0; i < n; ++i) - { - CV_Assert( src[i].size() == size ); - CV_Assert( src[i].depth() == depth ); - CV_Assert( src[i].channels() == 1 ); - } - - if (depth == CV_64F) - { - if (!deviceSupports(NATIVE_DOUBLE)) - CV_Error(cv::Error::StsUnsupportedFormat, "The device doesn't support double"); - } - - if (n == 1) - { - src[0].copyTo(_dst, stream); - } - else - { - _dst.create(size, CV_MAKE_TYPE(depth, (int)n)); - GpuMat dst = _dst.getGpuMat(); - - PtrStepSzb src_as_devmem[4]; - for(size_t i = 0; i < n; ++i) - src_as_devmem[i] = src[i]; - - PtrStepSzb dst_as_devmem(dst); - cv::cuda::device::split_merge::merge(src_as_devmem, dst_as_devmem, (int)n, CV_ELEM_SIZE(depth), StreamAccessor::getStream(stream)); - } - } - - void split_caller(const GpuMat& src, GpuMat* dst, Stream& stream) - { - CV_Assert( dst != 0 ); - - const int depth = src.depth(); - const int num_channels = src.channels(); - - CV_Assert( num_channels <= 4 ); - - if (depth == CV_64F) - { - if (!deviceSupports(NATIVE_DOUBLE)) - CV_Error(cv::Error::StsUnsupportedFormat, "The device doesn't support double"); - } - - if (num_channels == 1) - { - src.copyTo(dst[0], stream); - return; - } - - for (int i = 0; i < num_channels; ++i) - dst[i].create(src.size(), depth); - - PtrStepSzb dst_as_devmem[4]; - for (int i = 0; i < num_channels; ++i) - dst_as_devmem[i] = dst[i]; - - PtrStepSzb src_as_devmem(src); - cv::cuda::device::split_merge::split(src_as_devmem, dst_as_devmem, num_channels, src.elemSize1(), StreamAccessor::getStream(stream)); - } -} - -void cv::cuda::merge(const GpuMat* src, size_t n, OutputArray dst, Stream& stream) -{ - merge_caller(src, n, dst, stream); -} - - -void cv::cuda::merge(const std::vector& src, OutputArray dst, Stream& stream) -{ - merge_caller(&src[0], src.size(), dst, stream); -} - -void cv::cuda::split(InputArray _src, GpuMat* dst, Stream& stream) -{ - GpuMat src = _src.getGpuMat(); - split_caller(src, dst, stream); -} - -void cv::cuda::split(InputArray _src, std::vector& dst, Stream& stream) -{ - GpuMat src = _src.getGpuMat(); - dst.resize(src.channels()); - if(src.channels() > 0) - split_caller(src, &dst[0], stream); -} - -//////////////////////////////////////////////////////////////////////// -// transpose - -namespace arithm -{ - template void transpose(PtrStepSz src, PtrStepSz dst, cudaStream_t stream); -} - -void cv::cuda::transpose(InputArray _src, OutputArray _dst, Stream& _stream) -{ - GpuMat src = _src.getGpuMat(); - - CV_Assert( src.elemSize() == 1 || src.elemSize() == 4 || src.elemSize() == 8 ); - - _dst.create( src.cols, src.rows, src.type() ); - GpuMat dst = _dst.getGpuMat(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - if (src.elemSize() == 1) - { - NppStreamHandler h(stream); - - NppiSize sz; - sz.width = src.cols; - sz.height = src.rows; - - nppSafeCall( nppiTranspose_8u_C1R(src.ptr(), static_cast(src.step), - dst.ptr(), static_cast(dst.step), sz) ); - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } - else if (src.elemSize() == 4) - { - arithm::transpose(src, dst, stream); - } - else // if (src.elemSize() == 8) - { - if (!deviceSupports(NATIVE_DOUBLE)) - CV_Error(cv::Error::StsUnsupportedFormat, "The device doesn't support double"); - - arithm::transpose(src, dst, stream); - } -} - //////////////////////////////////////////////////////////////////////// // flip @@ -287,326 +130,4 @@ void cv::cuda::flip(InputArray _src, OutputArray _dst, int flipCode, Stream& str funcs[src.depth()][src.channels() - 1](src, dst, flipCode, StreamAccessor::getStream(stream)); } -//////////////////////////////////////////////////////////////////////// -// LUT - -#if (CUDA_VERSION >= 5000) - -namespace -{ - class LookUpTableImpl : public LookUpTable - { - public: - LookUpTableImpl(InputArray lut); - - void transform(InputArray src, OutputArray dst, Stream& stream = Stream::Null()); - - private: - int lut_cn; - - int nValues3[3]; - const Npp32s* pValues3[3]; - const Npp32s* pLevels3[3]; - - GpuMat d_pLevels; - GpuMat d_nppLut; - GpuMat d_nppLut3[3]; - }; - - LookUpTableImpl::LookUpTableImpl(InputArray _lut) - { - nValues3[0] = nValues3[1] = nValues3[2] = 256; - - Npp32s pLevels[256]; - for (int i = 0; i < 256; ++i) - pLevels[i] = i; - - d_pLevels.upload(Mat(1, 256, CV_32S, pLevels)); - pLevels3[0] = pLevels3[1] = pLevels3[2] = d_pLevels.ptr(); - - GpuMat lut; - if (_lut.kind() == _InputArray::GPU_MAT) - { - lut = _lut.getGpuMat(); - } - else - { - Mat hLut = _lut.getMat(); - CV_Assert( hLut.total() == 256 && hLut.isContinuous() ); - lut.upload(Mat(1, 256, hLut.type(), hLut.data)); - } - - lut_cn = lut.channels(); - - CV_Assert( lut.depth() == CV_8U ); - CV_Assert( lut.rows == 1 && lut.cols == 256 ); - - lut.convertTo(d_nppLut, CV_32S); - - if (lut_cn == 1) - { - pValues3[0] = pValues3[1] = pValues3[2] = d_nppLut.ptr(); - } - else - { - cuda::split(d_nppLut, d_nppLut3); - - pValues3[0] = d_nppLut3[0].ptr(); - pValues3[1] = d_nppLut3[1].ptr(); - pValues3[2] = d_nppLut3[2].ptr(); - } - } - - void LookUpTableImpl::transform(InputArray _src, OutputArray _dst, Stream& _stream) - { - GpuMat src = _src.getGpuMat(); - - const int cn = src.channels(); - - CV_Assert( src.type() == CV_8UC1 || src.type() == CV_8UC3 ); - CV_Assert( lut_cn == 1 || lut_cn == cn ); - - _dst.create(src.size(), src.type()); - GpuMat dst = _dst.getGpuMat(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - NppStreamHandler h(stream); - - NppiSize sz; - sz.height = src.rows; - sz.width = src.cols; - - if (src.type() == CV_8UC1) - { - nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr(), static_cast(src.step), - dst.ptr(), static_cast(dst.step), sz, d_nppLut.ptr(), d_pLevels.ptr(), 256) ); - } - else - { - nppSafeCall( nppiLUT_Linear_8u_C3R(src.ptr(), static_cast(src.step), - dst.ptr(), static_cast(dst.step), sz, pValues3, pLevels3, nValues3) ); - } - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } -} - -#else // (CUDA_VERSION >= 5000) - -namespace -{ - class LookUpTableImpl : public LookUpTable - { - public: - LookUpTableImpl(InputArray lut); - - void transform(InputArray src, OutputArray dst, Stream& stream = Stream::Null()); - - private: - int lut_cn; - - Npp32s pLevels[256]; - int nValues3[3]; - const Npp32s* pValues3[3]; - const Npp32s* pLevels3[3]; - - Mat nppLut; - Mat nppLut3[3]; - }; - - LookUpTableImpl::LookUpTableImpl(InputArray _lut) - { - nValues3[0] = nValues3[1] = nValues3[2] = 256; - - for (int i = 0; i < 256; ++i) - pLevels[i] = i; - pLevels3[0] = pLevels3[1] = pLevels3[2] = pLevels; - - Mat lut; - if (_lut.kind() == _InputArray::GPU_MAT) - { - lut = Mat(_lut.getGpuMat()); - } - else - { - Mat hLut = _lut.getMat(); - CV_Assert( hLut.total() == 256 && hLut.isContinuous() ); - lut = hLut; - } - - lut_cn = lut.channels(); - - CV_Assert( lut.depth() == CV_8U ); - CV_Assert( lut.rows == 1 && lut.cols == 256 ); - - lut.convertTo(nppLut, CV_32S); - - if (lut_cn == 1) - { - pValues3[0] = pValues3[1] = pValues3[2] = nppLut.ptr(); - } - else - { - cv::split(nppLut, nppLut3); - - pValues3[0] = nppLut3[0].ptr(); - pValues3[1] = nppLut3[1].ptr(); - pValues3[2] = nppLut3[2].ptr(); - } - } - - void LookUpTableImpl::transform(InputArray _src, OutputArray _dst, Stream& _stream) - { - GpuMat src = _src.getGpuMat(); - - const int cn = src.channels(); - - CV_Assert( src.type() == CV_8UC1 || src.type() == CV_8UC3 ); - CV_Assert( lut_cn == 1 || lut_cn == cn ); - - _dst.create(src.size(), src.type()); - GpuMat dst = _dst.getGpuMat(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - NppStreamHandler h(stream); - - NppiSize sz; - sz.height = src.rows; - sz.width = src.cols; - - if (src.type() == CV_8UC1) - { - nppSafeCall( nppiLUT_Linear_8u_C1R(src.ptr(), static_cast(src.step), - dst.ptr(), static_cast(dst.step), sz, nppLut.ptr(), pLevels, 256) ); - } - else - { - nppSafeCall( nppiLUT_Linear_8u_C3R(src.ptr(), static_cast(src.step), - dst.ptr(), static_cast(dst.step), sz, pValues3, pLevels3, nValues3) ); - } - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } -} - -#endif // (CUDA_VERSION >= 5000) - -Ptr cv::cuda::createLookUpTable(InputArray lut) -{ - return makePtr(lut); -} - -//////////////////////////////////////////////////////////////////////// -// copyMakeBorder - -namespace cv { namespace cuda { namespace device -{ - namespace imgproc - { - template void copyMakeBorder_gpu(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderMode, const T* borderValue, cudaStream_t stream); - } -}}} - -namespace -{ - template void copyMakeBorder_caller(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream) - { - using namespace ::cv::cuda::device::imgproc; - - Scalar_ val(saturate_cast(value[0]), saturate_cast(value[1]), saturate_cast(value[2]), saturate_cast(value[3])); - - copyMakeBorder_gpu(src, dst, top, left, borderType, val.val, stream); - } -} - -#if defined __GNUC__ && __GNUC__ > 2 && __GNUC_MINOR__ > 4 -typedef Npp32s __attribute__((__may_alias__)) Npp32s_a; -#else -typedef Npp32s Npp32s_a; -#endif - -void cv::cuda::copyMakeBorder(InputArray _src, OutputArray _dst, int top, int bottom, int left, int right, int borderType, Scalar value, Stream& _stream) -{ - GpuMat src = _src.getGpuMat(); - - CV_Assert( src.depth() <= CV_32F && src.channels() <= 4 ); - CV_Assert( borderType == BORDER_REFLECT_101 || borderType == BORDER_REPLICATE || borderType == BORDER_CONSTANT || borderType == BORDER_REFLECT || borderType == BORDER_WRAP ); - - _dst.create(src.rows + top + bottom, src.cols + left + right, src.type()); - GpuMat dst = _dst.getGpuMat(); - - cudaStream_t stream = StreamAccessor::getStream(_stream); - - if (borderType == BORDER_CONSTANT && (src.type() == CV_8UC1 || src.type() == CV_8UC4 || src.type() == CV_32SC1 || src.type() == CV_32FC1)) - { - NppiSize srcsz; - srcsz.width = src.cols; - srcsz.height = src.rows; - - NppiSize dstsz; - dstsz.width = dst.cols; - dstsz.height = dst.rows; - - NppStreamHandler h(stream); - - switch (src.type()) - { - case CV_8UC1: - { - Npp8u nVal = saturate_cast(value[0]); - nppSafeCall( nppiCopyConstBorder_8u_C1R(src.ptr(), static_cast(src.step), srcsz, - dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); - break; - } - case CV_8UC4: - { - Npp8u nVal[] = {saturate_cast(value[0]), saturate_cast(value[1]), saturate_cast(value[2]), saturate_cast(value[3])}; - nppSafeCall( nppiCopyConstBorder_8u_C4R(src.ptr(), static_cast(src.step), srcsz, - dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); - break; - } - case CV_32SC1: - { - Npp32s nVal = saturate_cast(value[0]); - nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr(), static_cast(src.step), srcsz, - dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); - break; - } - case CV_32FC1: - { - Npp32f val = saturate_cast(value[0]); - Npp32s nVal = *(reinterpret_cast(&val)); - nppSafeCall( nppiCopyConstBorder_32s_C1R(src.ptr(), static_cast(src.step), srcsz, - dst.ptr(), static_cast(dst.step), dstsz, top, left, nVal) ); - break; - } - } - - if (stream == 0) - cudaSafeCall( cudaDeviceSynchronize() ); - } - else - { - typedef void (*caller_t)(const PtrStepSzb& src, const PtrStepSzb& dst, int top, int left, int borderType, const Scalar& value, cudaStream_t stream); - static const caller_t callers[6][4] = - { - { copyMakeBorder_caller , copyMakeBorder_caller , copyMakeBorder_caller , copyMakeBorder_caller}, - {0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/ , 0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/}, - { copyMakeBorder_caller , 0/*copyMakeBorder_caller*/, copyMakeBorder_caller , copyMakeBorder_caller}, - { copyMakeBorder_caller , 0/*copyMakeBorder_caller*/ , copyMakeBorder_caller , copyMakeBorder_caller}, - {0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/ , 0/*copyMakeBorder_caller*/, 0/*copyMakeBorder_caller*/}, - { copyMakeBorder_caller , 0/*copyMakeBorder_caller*/ , copyMakeBorder_caller , copyMakeBorder_caller} - }; - - caller_t func = callers[src.depth()][src.channels() - 1]; - CV_Assert(func != 0); - - func(src, dst, top, left, borderType, value, stream); - } -} - #endif /* !defined (HAVE_CUDA) */ diff --git a/modules/cudaarithm/src/cuda/absdiff_mat.cu b/modules/cudaarithm/src/cuda/absdiff_mat.cu index ffdcbfa380..ec04f12284 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 <= CV_64F ); + + 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..3ffd0661b7 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 <= CV_64F ); + + funcs[depth](src, val[0], dst, stream); +} + +#endif diff --git a/modules/cudaarithm/src/cuda/add_mat.cu b/modules/cudaarithm/src/cuda/add_mat.cu index 1270438208..4166cc104e 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 <= CV_64F && ddepth <= CV_64F ); + + 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..92838a2a57 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 <= CV_64F && ddepth <= CV_64F && 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/cuda/add_weighted.cu b/modules/cudaarithm/src/cuda/add_weighted.cu index aa305d993e..d5c00f6072 100644 --- a/modules/cudaarithm/src/cuda/add_weighted.cu +++ b/modules/cudaarithm/src/cuda/add_weighted.cu @@ -40,325 +40,553 @@ // //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" +#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/cudaarithm.hpp" +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +namespace { - template struct UseDouble_ + template struct AddWeightedOp : binary_function { - enum {value = 0}; - }; - template <> struct UseDouble_ - { - enum {value = 1}; - }; - template struct UseDouble - { - enum {value = (UseDouble_::value || UseDouble_::value || UseDouble_::value)}; - }; - - template struct AddWeighted_; - template struct AddWeighted_ : binary_function - { - float alpha; - float beta; - float gamma; - - __host__ AddWeighted_(double alpha_, double beta_, double gamma_) : alpha(static_cast(alpha_)), beta(static_cast(beta_)), gamma(static_cast(gamma_)) {} + S alpha; + S beta; + S gamma; __device__ __forceinline__ D operator ()(T1 a, T2 b) const { return saturate_cast(a * alpha + b * beta + gamma); } }; - template struct AddWeighted_ : binary_function - { - double alpha; - double beta; - double gamma; - __host__ AddWeighted_(double alpha_, double beta_, double gamma_) : alpha(alpha_), beta(beta_), gamma(gamma_) {} - - __device__ __forceinline__ D operator ()(T1 a, T2 b) const - { - return saturate_cast(a * alpha + b * beta + gamma); - } - }; - template struct AddWeighted : AddWeighted_::value> + template struct TransformPolicy : DefaultTransformPolicy { - AddWeighted(double alpha_, double beta_, double gamma_) : AddWeighted_::value>(alpha_, beta_, gamma_) {} }; + template <> struct TransformPolicy : DefaultTransformPolicy + { + enum { + shift = 1 + }; + }; + + template + void addWeightedImpl(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst, Stream& stream) + { + typedef typename LargerType::type larger_type1; + typedef typename LargerType::type larger_type2; + typedef typename LargerType::type scalar_type; + + AddWeightedOp op; + op.alpha = static_cast(alpha); + op.beta = static_cast(beta); + op.gamma = static_cast(gamma); + + gridTransformBinary_< TransformPolicy >(globPtr(src1), globPtr(src2), globPtr(dst), op, stream); + } } -namespace cv { namespace cuda { namespace device +void cv::cuda::addWeighted(InputArray _src1, double alpha, InputArray _src2, double beta, double gamma, OutputArray _dst, int ddepth, Stream& stream) { - template struct AddWeightedTraits : DefaultTransformFunctorTraits< arithm::AddWeighted > - { - }; - template struct AddWeightedTraits : arithm::ArithmFuncTraits + typedef void (*func_t)(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst, Stream& stream); + static const func_t funcs[7][7][7] = { + { + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + } + }, + { + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + } + }, + { + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + } + }, + { + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + } + }, + { + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + } + }, + { + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + } + }, + { + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/, + 0/*addWeightedImpl*/ + }, + { + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl, + addWeightedImpl + } + } }; - template struct TransformFunctorTraits< arithm::AddWeighted > : AddWeightedTraits - { - }; -}}} + GpuMat src1 = _src1.getGpuMat(); + GpuMat src2 = _src2.getGpuMat(); -namespace arithm -{ - template - void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream) - { - AddWeighted op(alpha, beta, gamma); + int sdepth1 = src1.depth(); + int sdepth2 = src2.depth(); - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, op, WithOutMask(), stream); + ddepth = ddepth >= 0 ? CV_MAT_DEPTH(ddepth) : std::max(sdepth1, sdepth2); + const int cn = src1.channels(); + + CV_DbgAssert( src2.size() == src1.size() && src2.channels() == cn ); + CV_DbgAssert( sdepth1 <= CV_64F && sdepth2 <= CV_64F && ddepth <= CV_64F ); + + _dst.create(src1.size(), CV_MAKE_TYPE(ddepth, cn)); + GpuMat dst = _dst.getGpuMat(); + + GpuMat src1_ = src1.reshape(1); + GpuMat src2_ = src2.reshape(1); + GpuMat dst_ = dst.reshape(1); + + if (sdepth1 > sdepth2) + { + src1_.swap(src2_); + std::swap(alpha, beta); + std::swap(sdepth1, sdepth2); } - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); + const func_t func = funcs[sdepth1][sdepth2][ddepth]; - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); + if (!func) + CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - - - - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - template void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); + func(src1_, alpha, src2_, beta, gamma, dst_, stream); } -#endif /* CUDA_DISABLER */ +#endif diff --git a/modules/cudaarithm/src/cuda/arithm_func_traits.hpp b/modules/cudaarithm/src/cuda/arithm_func_traits.hpp deleted file mode 100644 index 051fd9ff08..0000000000 --- a/modules/cudaarithm/src/cuda/arithm_func_traits.hpp +++ /dev/null @@ -1,145 +0,0 @@ -/*M/////////////////////////////////////////////////////////////////////////////////////// -// -// IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING. -// -// By downloading, copying, installing or using the software you agree to this license. -// If you do not agree to this license, do not download, install, -// copy or use the software. -// -// -// License Agreement -// For Open Source Computer Vision Library -// -// Copyright (C) 2000-2008, Intel Corporation, all rights reserved. -// Copyright (C) 2009, Willow Garage Inc., all rights reserved. -// Third party copyrights are property of their respective owners. -// -// Redistribution and use in source and binary forms, with or without modification, -// are permitted provided that the following conditions are met: -// -// * Redistribution's of source code must retain the above copyright notice, -// this list of conditions and the following disclaimer. -// -// * Redistribution's in binary form must reproduce the above copyright notice, -// this list of conditions and the following disclaimer in the documentation -// and/or other materials provided with the distribution. -// -// * The name of the copyright holders may not be used to endorse or promote products -// derived from this software without specific prior written permission. -// -// This software is provided by the copyright holders and contributors "as is" and -// any express or implied warranties, including, but not limited to, the implied -// warranties of merchantability and fitness for a particular purpose are disclaimed. -// In no event shall the Intel Corporation or contributors be liable for any direct, -// indirect, incidental, special, exemplary, or consequential damages -// (including, but not limited to, procurement of substitute goods or services; -// loss of use, data, or profits; or business interruption) however caused -// and on any theory of liability, whether in contract, strict liability, -// or tort (including negligence or otherwise) arising in any way out of -// the use of this software, even if advised of the possibility of such damage. -// -//M*/ - -#ifndef __ARITHM_FUNC_TRAITS_HPP__ -#define __ARITHM_FUNC_TRAITS_HPP__ - -#include - -namespace arithm -{ - template struct ArithmFuncTraits - { - enum { simple_block_dim_x = 32 }; - enum { simple_block_dim_y = 8 }; - - enum { smart_block_dim_x = 32 }; - enum { smart_block_dim_y = 8 }; - enum { smart_shift = 1 }; - }; - - template <> struct ArithmFuncTraits<1, 1> - { - enum { simple_block_dim_x = 32 }; - enum { simple_block_dim_y = 8 }; - - enum { smart_block_dim_x = 32 }; - enum { smart_block_dim_y = 8 }; - enum { smart_shift = 4 }; - }; - template <> struct ArithmFuncTraits<1, 2> - { - enum { simple_block_dim_x = 32 }; - enum { simple_block_dim_y = 8 }; - - enum { smart_block_dim_x = 32 }; - enum { smart_block_dim_y = 8 }; - enum { smart_shift = 4 }; - }; - template <> struct ArithmFuncTraits<1, 4> - { - enum { simple_block_dim_x = 32 }; - enum { simple_block_dim_y = 8 }; - - enum { smart_block_dim_x = 32 }; - enum { smart_block_dim_y = 8 }; - enum { smart_shift = 4 }; - }; - - template <> struct ArithmFuncTraits<2, 1> - { - enum { simple_block_dim_x = 32 }; - enum { simple_block_dim_y = 8 }; - - enum { smart_block_dim_x = 32 }; - enum { smart_block_dim_y = 8 }; - enum { smart_shift = 4 }; - }; - template <> struct ArithmFuncTraits<2, 2> - { - enum { simple_block_dim_x = 32 }; - enum { simple_block_dim_y = 8 }; - - enum { smart_block_dim_x = 32 }; - enum { smart_block_dim_y = 8 }; - enum { smart_shift = 4 }; - }; - template <> struct ArithmFuncTraits<2, 4> - { - enum { simple_block_dim_x = 32 }; - enum { simple_block_dim_y = 8 }; - - enum { smart_block_dim_x = 32 }; - enum { smart_block_dim_y = 8 }; - enum { smart_shift = 4 }; - }; - - template <> struct ArithmFuncTraits<4, 1> - { - enum { simple_block_dim_x = 32 }; - enum { simple_block_dim_y = 8 }; - - enum { smart_block_dim_x = 32 }; - enum { smart_block_dim_y = 8 }; - enum { smart_shift = 4 }; - }; - template <> struct ArithmFuncTraits<4, 2> - { - enum { simple_block_dim_x = 32 }; - enum { simple_block_dim_y = 8 }; - - enum { smart_block_dim_x = 32 }; - enum { smart_block_dim_y = 8 }; - enum { smart_shift = 4 }; - }; - template <> struct ArithmFuncTraits<4, 4> - { - enum { simple_block_dim_x = 32 }; - enum { simple_block_dim_y = 8 }; - - enum { smart_block_dim_x = 32 }; - enum { smart_block_dim_y = 8 }; - enum { smart_shift = 4 }; - }; -} - -#endif // __ARITHM_FUNC_TRAITS_HPP__ diff --git a/modules/cudaarithm/src/cuda/bitwise_mat.cu b/modules/cudaarithm/src/cuda/bitwise_mat.cu index 7a90cf3fc3..b2bf288be7 100644 --- a/modules/cudaarithm/src/cuda/bitwise_mat.cu +++ b/modules/cudaarithm/src/cuda/bitwise_mat.cu @@ -40,87 +40,187 @@ // //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 cv { namespace cuda { namespace device +#include "opencv2/cudaarithm.hpp" +#include "opencv2/cudev.hpp" + +using namespace cv::cudev; + +void bitMat(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, double, Stream& stream, int op); + +////////////////////////////////////////////////////////////////////////////// +/// bitwise_not + +void cv::cuda::bitwise_not(InputArray _src, OutputArray _dst, InputArray _mask, Stream& stream) { - template struct TransformFunctorTraits< bit_not > : arithm::ArithmFuncTraits - { - }; + GpuMat src = _src.getGpuMat(); + GpuMat mask = _mask.getGpuMat(); - template struct TransformFunctorTraits< bit_and > : arithm::ArithmFuncTraits - { - }; + const int depth = src.depth(); - template struct TransformFunctorTraits< bit_or > : arithm::ArithmFuncTraits - { - }; + CV_DbgAssert( depth <= CV_32F ); + CV_DbgAssert( mask.empty() || (mask.type() == CV_8UC1 && mask.size() == src.size()) ); - template struct TransformFunctorTraits< bit_xor > : arithm::ArithmFuncTraits - { - }; -}}} + _dst.create(src.size(), src.type()); + GpuMat dst = _dst.getGpuMat(); -namespace arithm -{ - template void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) + if (mask.empty()) { - if (mask.data) - device::transform((PtrStepSz) src, (PtrStepSz) dst, bit_not(), mask, stream); + const int bcols = (int) (src.cols * src.elemSize()); + + if ((bcols & 3) == 0) + { + const int vcols = bcols >> 2; + + GlobPtrSz vsrc = globPtr((uint*) src.data, src.step, src.rows, vcols); + GlobPtrSz vdst = globPtr((uint*) dst.data, dst.step, src.rows, vcols); + + gridTransformUnary(vsrc, vdst, bit_not(), stream); + } + else if ((bcols & 1) == 0) + { + const int vcols = bcols >> 1; + + GlobPtrSz vsrc = globPtr((ushort*) src.data, src.step, src.rows, vcols); + GlobPtrSz vdst = globPtr((ushort*) dst.data, dst.step, src.rows, vcols); + + gridTransformUnary(vsrc, vdst, bit_not(), stream); + } else - device::transform((PtrStepSz) src, (PtrStepSz) dst, bit_not(), WithOutMask(), stream); - } + { + GlobPtrSz vsrc = globPtr((uchar*) src.data, src.step, src.rows, bcols); + GlobPtrSz vdst = globPtr((uchar*) dst.data, dst.step, src.rows, bcols); - template void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) + gridTransformUnary(vsrc, vdst, bit_not(), stream); + } + } + else { - if (mask.data) - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, bit_and(), mask, stream); + if (depth == CV_32F || depth == CV_32S) + { + GlobPtrSz vsrc = globPtr((uint*) src.data, src.step, src.rows, src.cols * src.channels()); + GlobPtrSz vdst = globPtr((uint*) dst.data, dst.step, src.rows, src.cols * src.channels()); + + gridTransformUnary(vsrc, vdst, bit_not(), singleMaskChannels(globPtr(mask), src.channels()), stream); + } + else if (depth == CV_16S || depth == CV_16U) + { + GlobPtrSz vsrc = globPtr((ushort*) src.data, src.step, src.rows, src.cols * src.channels()); + GlobPtrSz vdst = globPtr((ushort*) dst.data, dst.step, src.rows, src.cols * src.channels()); + + gridTransformUnary(vsrc, vdst, bit_not(), singleMaskChannels(globPtr(mask), src.channels()), stream); + } else - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, bit_and(), WithOutMask(), stream); + { + GlobPtrSz vsrc = globPtr((uchar*) src.data, src.step, src.rows, src.cols * src.channels()); + GlobPtrSz vdst = globPtr((uchar*) dst.data, dst.step, src.rows, src.cols * src.channels()); + + gridTransformUnary(vsrc, vdst, bit_not(), singleMaskChannels(globPtr(mask), src.channels()), stream); + } } - - template void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) - { - if (mask.data) - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, bit_or(), mask, stream); - else - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, bit_or(), WithOutMask(), stream); - } - - template void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream) - { - if (mask.data) - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, bit_xor(), mask, stream); - else - device::transform((PtrStepSz) src1, (PtrStepSz) src2, (PtrStepSz) dst, bit_xor(), WithOutMask(), stream); - } - - template void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void bitMatNot(PtrStepSzb src, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - template void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void bitMatAnd(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - template void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void bitMatOr(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - - template void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); - template void bitMatXor(PtrStepSzb src1, PtrStepSzb src2, PtrStepSzb dst, PtrStepb mask, cudaStream_t stream); } -#endif // CUDA_DISABLER +////////////////////////////////////////////////////////////////////////////// +/// Binary bitwise logical operations + +namespace +{ + template