From e83be009a39b8b2c147a39de558214ca71bdbb0f Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 26 Aug 2013 10:30:04 +0400 Subject: [PATCH] used new device layer for cv::gpu::addWeighted --- modules/cudaarithm/src/cuda/add_weighted.cu | 814 +++++++++++------- modules/cudaarithm/src/element_operations.cpp | 511 ----------- 2 files changed, 521 insertions(+), 804 deletions(-) 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/element_operations.cpp b/modules/cudaarithm/src/element_operations.cpp index f05004a27a..11d6b87447 100644 --- a/modules/cudaarithm/src/element_operations.cpp +++ b/modules/cudaarithm/src/element_operations.cpp @@ -449,517 +449,6 @@ void cv::cuda::max(InputArray src1, InputArray src2, OutputArray dst, Stream& st arithm_op(src1, src2, dst, noArray(), 1.0, -1, stream, minMaxMat, minMaxScalar, MAX_OP); } -//////////////////////////////////////////////////////////////////////// -// addWeighted - -namespace arithm -{ - template - void addWeighted(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); -} - -void cv::cuda::addWeighted(InputArray _src1, double alpha, InputArray _src2, double beta, double gamma, OutputArray _dst, int ddepth, Stream& stream) -{ - typedef void (*func_t)(PtrStepSzb src1, double alpha, PtrStepSzb src2, double beta, double gamma, PtrStepSzb dst, cudaStream_t stream); - static const func_t funcs[7][7][7] = - { - { - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - } - }, - { - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - } - }, - { - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - } - }, - { - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - } - }, - { - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - } - }, - { - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - } - }, - { - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/, - 0/*arithm::addWeighted*/ - }, - { - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted, - arithm::addWeighted - } - } - }; - - GpuMat src1 = _src1.getGpuMat(); - GpuMat src2 = _src2.getGpuMat(); - - int sdepth1 = src1.depth(); - int sdepth2 = src2.depth(); - ddepth = ddepth >= 0 ? CV_MAT_DEPTH(ddepth) : std::max(sdepth1, sdepth2); - const int cn = src1.channels(); - - CV_Assert( src2.size() == src1.size() && src2.channels() == cn ); - CV_Assert( sdepth1 <= CV_64F && sdepth2 <= CV_64F && ddepth <= CV_64F ); - - if (sdepth1 == CV_64F || sdepth2 == CV_64F || ddepth == CV_64F) - { - if (!deviceSupports(NATIVE_DOUBLE)) - CV_Error(cv::Error::StsUnsupportedFormat, "The device doesn't support double"); - } - - _dst.create(src1.size(), CV_MAKE_TYPE(ddepth, cn)); - GpuMat dst = _dst.getGpuMat(); - - PtrStepSzb src1_(src1.rows, src1.cols * cn, src1.data, src1.step); - PtrStepSzb src2_(src1.rows, src1.cols * cn, src2.data, src2.step); - PtrStepSzb dst_(src1.rows, src1.cols * cn, dst.data, dst.step); - - if (sdepth1 > sdepth2) - { - std::swap(src1_.data, src2_.data); - std::swap(src1_.step, src2_.step); - std::swap(alpha, beta); - std::swap(sdepth1, sdepth2); - } - - const func_t func = funcs[sdepth1][sdepth2][ddepth]; - - if (!func) - CV_Error(cv::Error::StsUnsupportedFormat, "Unsupported combination of source and destination types"); - - func(src1_, alpha, src2_, beta, gamma, dst_, StreamAccessor::getStream(stream)); -} - //////////////////////////////////////////////////////////////////////// // threshold