From 0aaaad1ea87eb719044e3c7b125965f934e7a636 Mon Sep 17 00:00:00 2001
From: Vladislav Vinogradov <no@email>
Date: Wed, 21 Sep 2011 08:58:54 +0000
Subject: [PATCH] implemented gpu::addWeighted

---
 modules/gpu/include/opencv2/gpu/gpu.hpp    |   4 +
 modules/gpu/perf/perf_arithm.cpp           |  31 ++
 modules/gpu/src/cuda/element_operations.cu | 322 +++++++++++++
 modules/gpu/src/element_operations.cpp     | 506 ++++++++++++++++++++-
 modules/gpu/test/test_arithm.cpp           |  96 +++-
 5 files changed, 943 insertions(+), 16 deletions(-)

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index dbb110741..7377f5257 100644
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -593,6 +593,10 @@ namespace cv
         //! computes per-element maximum of array and scalar (dst = max(src1, src2))
         CV_EXPORTS void max(const GpuMat& src1, double src2, GpuMat& dst, Stream& stream = Stream::Null());
 
+        //! computes the weighted sum of two arrays
+        CV_EXPORTS void addWeighted(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst, 
+            int dtype = -1, Stream& stream = Stream::Null());
+
 
         ////////////////////////////// Image processing //////////////////////////////
 
diff --git a/modules/gpu/perf/perf_arithm.cpp b/modules/gpu/perf/perf_arithm.cpp
index f5c4da3b0..2c915904c 100644
--- a/modules/gpu/perf/perf_arithm.cpp
+++ b/modules/gpu/perf/perf_arithm.cpp
@@ -685,3 +685,34 @@ PERF_TEST_P(DevInfo_Size_MatType, countNonZero, testing::Combine(testing::Values
 
     SANITY_CHECK(dst);
 }
+
+PERF_TEST_P(DevInfo_Size_MatType, addWeighted, testing::Combine(testing::ValuesIn(devices()), 
+                                                                testing::Values(GPU_TYPICAL_MAT_SIZES), 
+                                                                testing::Values(CV_8UC1, CV_16UC1, CV_32FC1)))
+{
+    DeviceInfo devInfo = std::tr1::get<0>(GetParam());
+    Size size = std::tr1::get<1>(GetParam());
+    int type = std::tr1::get<2>(GetParam());
+
+    setDevice(devInfo.deviceID());
+
+    Mat src1_host(size, type);
+    Mat src2_host(size, type);
+
+    declare.in(src1_host, src2_host, WARMUP_RNG);
+
+    GpuMat src1(src1_host);
+    GpuMat src2(src2_host);
+    GpuMat dst(size, type);
+
+    declare.time(0.5).iterations(100);
+
+    SIMPLE_TEST_CYCLE()
+    {
+        addWeighted(src1, 0.5, src2, 0.5, 0.0, dst);
+    }
+
+    Mat dst_host = dst;
+
+    SANITY_CHECK(dst_host);
+}
diff --git a/modules/gpu/src/cuda/element_operations.cu b/modules/gpu/src/cuda/element_operations.cu
index 71574b734..7548e901f 100644
--- a/modules/gpu/src/cuda/element_operations.cu
+++ b/modules/gpu/src/cuda/element_operations.cu
@@ -757,4 +757,326 @@ namespace cv { namespace gpu { namespace device
     }
 
     template void multiplyScalar_gpu<uchar, uchar>(const DevMem2D& src, float scale, const DevMem2D& dst, cudaStream_t stream);
+
+
+    //////////////////////////////////////////////////////////////////////////
+    // addWeighted
+
+    template <typename T1, typename T2, typename D> struct AddWeighted : binary_function<T1, T2, D>
+    {
+        __host__ __device__ __forceinline__ AddWeighted(double alpha_, double beta_, double gamma_) : alpha(alpha_), beta(beta_), gamma(gamma_) {}
+
+        __device__ __forceinline__ D operator ()(typename TypeTraits<T1>::ParameterType a, typename TypeTraits<T2>::ParameterType b) const
+        {
+            return saturate_cast<D>(alpha * a + beta * b + gamma);
+        }
+
+        const double alpha;
+        const double beta;
+        const double gamma;
+    };
+
+    template <> struct TransformFunctorTraits< AddWeighted<ushort, ushort, ushort> > : DefaultTransformFunctorTraits< AddWeighted<ushort, ushort, ushort> >
+    {
+        enum { smart_shift = 4 };
+    };
+    template <> struct TransformFunctorTraits< AddWeighted<ushort, ushort, short> > : DefaultTransformFunctorTraits< AddWeighted<ushort, ushort, short> >
+    {
+        enum { smart_shift = 4 };
+    };
+    template <> struct TransformFunctorTraits< AddWeighted<ushort, short, ushort> > : DefaultTransformFunctorTraits< AddWeighted<ushort, short, ushort> >
+    {
+        enum { smart_shift = 4 };
+    };
+    template <> struct TransformFunctorTraits< AddWeighted<ushort, short, short> > : DefaultTransformFunctorTraits< AddWeighted<ushort, short, short> >
+    {
+        enum { smart_shift = 4 };
+    };
+    template <> struct TransformFunctorTraits< AddWeighted<short, short, ushort> > : DefaultTransformFunctorTraits< AddWeighted<short, short, ushort> >
+    {
+        enum { smart_shift = 4 };
+    };
+    template <> struct TransformFunctorTraits< AddWeighted<short, short, short> > : DefaultTransformFunctorTraits< AddWeighted<short, short, short> >
+    {
+        enum { smart_shift = 4 };
+    };
+
+    template <> struct TransformFunctorTraits< AddWeighted<int, int, int> > : DefaultTransformFunctorTraits< AddWeighted<int, int, int> >
+    {
+        enum { smart_block_dim_y = 8 };
+        enum { smart_shift = 4 };
+    };
+    template <> struct TransformFunctorTraits< AddWeighted<int, int, float> > : DefaultTransformFunctorTraits< AddWeighted<int, int, float> >
+    {
+        enum { smart_block_dim_y = 8 };
+        enum { smart_shift = 4 };
+    };
+    template <> struct TransformFunctorTraits< AddWeighted<int, float, int> > : DefaultTransformFunctorTraits< AddWeighted<int, float, int> >
+    {
+        enum { smart_block_dim_y = 8 };
+        enum { smart_shift = 4 };
+    };
+    template <> struct TransformFunctorTraits< AddWeighted<int, float, float> > : DefaultTransformFunctorTraits< AddWeighted<int, float, float> >
+    {
+        enum { smart_block_dim_y = 8 };
+        enum { smart_shift = 4 };
+    };
+    template <> struct TransformFunctorTraits< AddWeighted<float, float, int> > : DefaultTransformFunctorTraits< AddWeighted<float, float, float> >
+    {
+        enum { smart_block_dim_y = 8 };
+        enum { smart_shift = 4 };
+    };
+    template <> struct TransformFunctorTraits< AddWeighted<float, float, float> > : DefaultTransformFunctorTraits< AddWeighted<float, float, float> >
+    {
+        enum { smart_block_dim_y = 8 };
+        enum { smart_shift = 4 };
+    };
+
+    template <typename T1, typename T2, typename D>
+    void addWeighted_gpu(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream)
+    {
+        cudaSafeCall( cudaSetDoubleForDevice(&alpha) );
+        cudaSafeCall( cudaSetDoubleForDevice(&beta) );
+        cudaSafeCall( cudaSetDoubleForDevice(&gamma) );
+
+        AddWeighted<T1, T2, D> op(alpha, beta, gamma);
+
+        transform(static_cast< DevMem2D_<T1> >(src1), static_cast< DevMem2D_<T2> >(src2), static_cast< DevMem2D_<D> >(dst), op, stream);
+    }
+
+    template void addWeighted_gpu<uchar, uchar, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, uchar, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, uchar, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, uchar, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, uchar, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, uchar, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, uchar, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<uchar, schar, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, schar, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, schar, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, schar, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, schar, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, schar, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, schar, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<uchar, ushort, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, ushort, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, ushort, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, ushort, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, ushort, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, ushort, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, ushort, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<uchar, short, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, short, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, short, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, short, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, short, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, short, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, short, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<uchar, int, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, int, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, int, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, int, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, int, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, int, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, int, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<uchar, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<uchar, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<uchar, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+
+
+    template void addWeighted_gpu<schar, schar, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, schar, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, schar, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, schar, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, schar, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, schar, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, schar, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<schar, ushort, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, ushort, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, ushort, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, ushort, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, ushort, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, ushort, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, ushort, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<schar, short, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, short, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, short, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, short, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, short, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, short, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, short, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<schar, int, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, int, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, int, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, int, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, int, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, int, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, int, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<schar, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<schar, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<schar, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+
+
+    template void addWeighted_gpu<ushort, ushort, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, ushort, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, ushort, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, ushort, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, ushort, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, ushort, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, ushort, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<ushort, short, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, short, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, short, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, short, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, short, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, short, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, short, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<ushort, int, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, int, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, int, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, int, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, int, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, int, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, int, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<ushort, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<ushort, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<ushort, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+
+
+    template void addWeighted_gpu<short, short, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, short, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, short, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, short, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, short, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, short, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, short, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<short, int, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, int, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, int, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, int, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, int, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, int, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, int, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<short, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<short, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<short, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    
+
+    template void addWeighted_gpu<int, int, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, int, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, int, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, int, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, int, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, int, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, int, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<int, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<int, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<int, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    
+
+    template void addWeighted_gpu<float, float, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, float, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, float, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, float, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, float, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, float, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, float, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    template void addWeighted_gpu<float, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<float, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    
+
+    template void addWeighted_gpu<double, double, uchar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<double, double, schar>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<double, double, ushort>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<double, double, short>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<double, double, int>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<double, double, float>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+    template void addWeighted_gpu<double, double, double>(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
 }}}
diff --git a/modules/gpu/src/element_operations.cpp b/modules/gpu/src/element_operations.cpp
index 68724fe66..4cf857ed4 100644
--- a/modules/gpu/src/element_operations.cpp
+++ b/modules/gpu/src/element_operations.cpp
@@ -67,7 +67,8 @@ void cv::gpu::min(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }
 void cv::gpu::max(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }
 void cv::gpu::max(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }
 double cv::gpu::threshold(const GpuMat&, GpuMat&, double, double, int, Stream&) {throw_nogpu(); return 0.0;}
-void cv::gpu::pow(const GpuMat&, double, GpuMat&, Stream&)  { throw_nogpu(); }
+void cv::gpu::pow(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }
+void cv::gpu::addWeighted(const GpuMat&, double, const GpuMat&, double, double, GpuMat&, int, Stream&) { throw_nogpu(); }
 
 #else
 
@@ -248,7 +249,7 @@ void cv::gpu::multiply(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream&
     {
         dst.create(src.size(), src.type());
 
-        device::multiplyScalar_gpu<uchar, uchar>(src.reshape(1), (float)(sc[0]), dst, StreamAccessor::getStream(stream));
+        device::multiplyScalar_gpu<unsigned char, unsigned char>(src.reshape(1), (float)(sc[0]), dst, StreamAccessor::getStream(stream));
     }
     else
     {
@@ -733,7 +734,7 @@ void cv::gpu::min(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& s
     typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream);
     static const func_t funcs[] = 
     {
-        min_caller<uchar>, min_caller<schar>, min_caller<ushort>, min_caller<short>, min_caller<int>, 
+        min_caller<unsigned char>, min_caller<signed char>, min_caller<unsigned short>, min_caller<short>, min_caller<int>, 
         min_caller<float>, min_caller<double>
     };
     funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));
@@ -746,7 +747,7 @@ void cv::gpu::min(const GpuMat& src1, double src2, GpuMat& dst, Stream& stream)
     typedef void (*func_t)(const GpuMat& src1, double src2, GpuMat& dst, cudaStream_t stream);
     static const func_t funcs[] = 
     {
-        min_caller<uchar>, min_caller<schar>, min_caller<ushort>, min_caller<short>, min_caller<int>, 
+        min_caller<unsigned char>, min_caller<signed char>, min_caller<unsigned short>, min_caller<short>, min_caller<int>, 
         min_caller<float>, min_caller<double>
     };
     funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));
@@ -761,7 +762,7 @@ void cv::gpu::max(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& s
     typedef void (*func_t)(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, cudaStream_t stream);
     static const func_t funcs[] = 
     {
-        max_caller<uchar>, max_caller<schar>, max_caller<ushort>, max_caller<short>, max_caller<int>, 
+        max_caller<unsigned char>, max_caller<signed char>, max_caller<unsigned short>, max_caller<short>, max_caller<int>, 
         max_caller<float>, max_caller<double>
     };
     funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));
@@ -775,7 +776,7 @@ void cv::gpu::max(const GpuMat& src1, double src2, GpuMat& dst, Stream& stream)
     typedef void (*func_t)(const GpuMat& src1, double src2, GpuMat& dst, cudaStream_t stream);
     static const func_t funcs[] = 
     {
-        max_caller<uchar>, max_caller<schar>, max_caller<ushort>, max_caller<short>, max_caller<int>, 
+        max_caller<unsigned char>, max_caller<signed char>, max_caller<unsigned short>, max_caller<short>, max_caller<int>, 
         max_caller<float>, max_caller<double>
     };
     funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));
@@ -875,4 +876,497 @@ void cv::gpu::pow(const GpuMat& src, double power, GpuMat& dst, Stream& stream)
     callers[src.depth()](src.reshape(1), (float)power, dst.reshape(1), StreamAccessor::getStream(stream));    
 }
 
+////////////////////////////////////////////////////////////////////////
+// addWeighted
+
+namespace cv { namespace gpu { namespace device
+{
+    template <typename T1, typename T2, typename D>
+    void addWeighted_gpu(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+}}}
+
+void cv::gpu::addWeighted(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst, int dtype, Stream& stream)
+{
+    CV_Assert(src1.size() == src2.size());
+    CV_Assert(src1.type() == src2.type() || (dtype >= 0 && src1.channels() == src2.channels()));
+
+    dtype = dtype >= 0 ? CV_MAKETYPE(dtype, src1.channels()) : src1.type();
+
+    dst.create(src1.size(), dtype);
+
+    const GpuMat* psrc1 = &src1;
+    const GpuMat* psrc2 = &src2;
+
+    if (src1.depth() > src2.depth())
+    {
+        std::swap(psrc1, psrc2);
+        std::swap(alpha, beta);
+    }
+
+    typedef void (*caller_t)(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);
+
+    using namespace cv::gpu::device;
+
+    static const caller_t callers[7][7][7] =
+    {
+        {
+            {
+                addWeighted_gpu<unsigned char, unsigned char, unsigned char >,
+                addWeighted_gpu<unsigned char, unsigned char, signed char >,
+                addWeighted_gpu<unsigned char, unsigned char, unsigned short>,
+                addWeighted_gpu<unsigned char, unsigned char, short >,
+                addWeighted_gpu<unsigned char, unsigned char, int   >,
+                addWeighted_gpu<unsigned char, unsigned char, float >,
+                addWeighted_gpu<unsigned char, unsigned char, double>
+            },
+            {
+                addWeighted_gpu<unsigned char, signed char, unsigned char >,
+                addWeighted_gpu<unsigned char, signed char, signed char >,
+                addWeighted_gpu<unsigned char, signed char, unsigned short>,
+                addWeighted_gpu<unsigned char, signed char, short >,
+                addWeighted_gpu<unsigned char, signed char, int   >,
+                addWeighted_gpu<unsigned char, signed char, float >,
+                addWeighted_gpu<unsigned char, signed char, double>
+            },
+            {
+                addWeighted_gpu<unsigned char, unsigned short, unsigned char >,
+                addWeighted_gpu<unsigned char, unsigned short, signed char >,
+                addWeighted_gpu<unsigned char, unsigned short, unsigned short>,
+                addWeighted_gpu<unsigned char, unsigned short, short >,
+                addWeighted_gpu<unsigned char, unsigned short, int   >,
+                addWeighted_gpu<unsigned char, unsigned short, float >,
+                addWeighted_gpu<unsigned char, unsigned short, double>
+            },
+            {
+                addWeighted_gpu<unsigned char, short, unsigned char >,
+                addWeighted_gpu<unsigned char, short, signed char >,
+                addWeighted_gpu<unsigned char, short, unsigned short>,
+                addWeighted_gpu<unsigned char, short, short >,
+                addWeighted_gpu<unsigned char, short, int   >,
+                addWeighted_gpu<unsigned char, short, float >,
+                addWeighted_gpu<unsigned char, short, double>
+            },
+            {
+                addWeighted_gpu<unsigned char, int, unsigned char >,
+                addWeighted_gpu<unsigned char, int, signed char >,
+                addWeighted_gpu<unsigned char, int, unsigned short>,
+                addWeighted_gpu<unsigned char, int, short >,
+                addWeighted_gpu<unsigned char, int, int   >,
+                addWeighted_gpu<unsigned char, int, float >,
+                addWeighted_gpu<unsigned char, int, double>
+            },
+            {
+                addWeighted_gpu<unsigned char, float, unsigned char >,
+                addWeighted_gpu<unsigned char, float, signed char >,
+                addWeighted_gpu<unsigned char, float, unsigned short>,
+                addWeighted_gpu<unsigned char, float, short >,
+                addWeighted_gpu<unsigned char, float, int   >,
+                addWeighted_gpu<unsigned char, float, float >,
+                addWeighted_gpu<unsigned char, float, double>
+            },
+            {
+                addWeighted_gpu<unsigned char, double, unsigned char >,
+                addWeighted_gpu<unsigned char, double, signed char >,
+                addWeighted_gpu<unsigned char, double, unsigned short>,
+                addWeighted_gpu<unsigned char, double, short >,
+                addWeighted_gpu<unsigned char, double, int   >,
+                addWeighted_gpu<unsigned char, double, float >,
+                addWeighted_gpu<unsigned char, double, double>
+            }
+        },
+        {
+            {
+                0/*addWeighted_gpu<signed char, unsigned char, unsigned char >*/,
+                0/*addWeighted_gpu<signed char, unsigned char, signed char >*/,
+                0/*addWeighted_gpu<signed char, unsigned char, unsigned short>*/,
+                0/*addWeighted_gpu<signed char, unsigned char, short >*/,
+                0/*addWeighted_gpu<signed char, unsigned char, int   >*/,
+                0/*addWeighted_gpu<signed char, unsigned char, float >*/,
+                0/*addWeighted_gpu<signed char, unsigned char, double>*/
+            },
+            {
+                addWeighted_gpu<signed char, signed char, unsigned char >,
+                addWeighted_gpu<signed char, signed char, signed char >,
+                addWeighted_gpu<signed char, signed char, unsigned short>,
+                addWeighted_gpu<signed char, signed char, short >,
+                addWeighted_gpu<signed char, signed char, int   >,
+                addWeighted_gpu<signed char, signed char, float >,
+                addWeighted_gpu<signed char, signed char, double>
+            },
+            {
+                addWeighted_gpu<signed char, unsigned short, unsigned char >,
+                addWeighted_gpu<signed char, unsigned short, signed char >,
+                addWeighted_gpu<signed char, unsigned short, unsigned short>,
+                addWeighted_gpu<signed char, unsigned short, short >,
+                addWeighted_gpu<signed char, unsigned short, int   >,
+                addWeighted_gpu<signed char, unsigned short, float >,
+                addWeighted_gpu<signed char, unsigned short, double>
+            },
+            {
+                addWeighted_gpu<signed char, short, unsigned char >,
+                addWeighted_gpu<signed char, short, signed char >,
+                addWeighted_gpu<signed char, short, unsigned short>,
+                addWeighted_gpu<signed char, short, short >,
+                addWeighted_gpu<signed char, short, int   >,
+                addWeighted_gpu<signed char, short, float >,
+                addWeighted_gpu<signed char, short, double>
+            },
+            {
+                addWeighted_gpu<signed char, int, unsigned char >,
+                addWeighted_gpu<signed char, int, signed char >,
+                addWeighted_gpu<signed char, int, unsigned short>,
+                addWeighted_gpu<signed char, int, short >,
+                addWeighted_gpu<signed char, int, int   >,
+                addWeighted_gpu<signed char, int, float >,
+                addWeighted_gpu<signed char, int, double>
+            },
+            {
+                addWeighted_gpu<signed char, float, unsigned char >,
+                addWeighted_gpu<signed char, float, signed char >,
+                addWeighted_gpu<signed char, float, unsigned short>,
+                addWeighted_gpu<signed char, float, short >,
+                addWeighted_gpu<signed char, float, int   >,
+                addWeighted_gpu<signed char, float, float >,
+                addWeighted_gpu<signed char, float, double>
+            },
+            {
+                addWeighted_gpu<signed char, double, unsigned char >,
+                addWeighted_gpu<signed char, double, signed char >,
+                addWeighted_gpu<signed char, double, unsigned short>,
+                addWeighted_gpu<signed char, double, short >,
+                addWeighted_gpu<signed char, double, int   >,
+                addWeighted_gpu<signed char, double, float >,
+                addWeighted_gpu<signed char, double, double>
+            }
+        },
+        {
+            {
+                0/*addWeighted_gpu<unsigned short, unsigned char, unsigned char >*/,
+                0/*addWeighted_gpu<unsigned short, unsigned char, signed char >*/,
+                0/*addWeighted_gpu<unsigned short, unsigned char, unsigned short>*/,
+                0/*addWeighted_gpu<unsigned short, unsigned char, short >*/,
+                0/*addWeighted_gpu<unsigned short, unsigned char, int   >*/,
+                0/*addWeighted_gpu<unsigned short, unsigned char, float >*/,
+                0/*addWeighted_gpu<unsigned short, unsigned char, double>*/
+            },
+            {
+                0/*addWeighted_gpu<unsigned short, signed char, unsigned char >*/,
+                0/*addWeighted_gpu<unsigned short, signed char, signed char >*/,
+                0/*addWeighted_gpu<unsigned short, signed char, unsigned short>*/,
+                0/*addWeighted_gpu<unsigned short, signed char, short >*/,
+                0/*addWeighted_gpu<unsigned short, signed char, int   >*/,
+                0/*addWeighted_gpu<unsigned short, signed char, float >*/,
+                0/*addWeighted_gpu<unsigned short, signed char, double>*/
+            },
+            {
+                addWeighted_gpu<unsigned short, unsigned short, unsigned char >,
+                addWeighted_gpu<unsigned short, unsigned short, signed char >,
+                addWeighted_gpu<unsigned short, unsigned short, unsigned short>,
+                addWeighted_gpu<unsigned short, unsigned short, short >,
+                addWeighted_gpu<unsigned short, unsigned short, int   >,
+                addWeighted_gpu<unsigned short, unsigned short, float >,
+                addWeighted_gpu<unsigned short, unsigned short, double>
+            },
+            {
+                addWeighted_gpu<unsigned short, short, unsigned char >,
+                addWeighted_gpu<unsigned short, short, signed char >,
+                addWeighted_gpu<unsigned short, short, unsigned short>,
+                addWeighted_gpu<unsigned short, short, short >,
+                addWeighted_gpu<unsigned short, short, int   >,
+                addWeighted_gpu<unsigned short, short, float >,
+                addWeighted_gpu<unsigned short, short, double>
+            },
+            {
+                addWeighted_gpu<unsigned short, int, unsigned char >,
+                addWeighted_gpu<unsigned short, int, signed char >,
+                addWeighted_gpu<unsigned short, int, unsigned short>,
+                addWeighted_gpu<unsigned short, int, short >,
+                addWeighted_gpu<unsigned short, int, int   >,
+                addWeighted_gpu<unsigned short, int, float >,
+                addWeighted_gpu<unsigned short, int, double>
+            },
+            {
+                addWeighted_gpu<unsigned short, float, unsigned char >,
+                addWeighted_gpu<unsigned short, float, signed char >,
+                addWeighted_gpu<unsigned short, float, unsigned short>,
+                addWeighted_gpu<unsigned short, float, short >,
+                addWeighted_gpu<unsigned short, float, int   >,
+                addWeighted_gpu<unsigned short, float, float >,
+                addWeighted_gpu<unsigned short, float, double>
+            },
+            {
+                addWeighted_gpu<unsigned short, double, unsigned char >,
+                addWeighted_gpu<unsigned short, double, signed char >,
+                addWeighted_gpu<unsigned short, double, unsigned short>,
+                addWeighted_gpu<unsigned short, double, short >,
+                addWeighted_gpu<unsigned short, double, int   >,
+                addWeighted_gpu<unsigned short, double, float >,
+                addWeighted_gpu<unsigned short, double, double>
+            }
+        },
+        {
+            {
+                0/*addWeighted_gpu<short, unsigned char, unsigned char >*/,
+                0/*addWeighted_gpu<short, unsigned char, signed char >*/,
+                0/*addWeighted_gpu<short, unsigned char, unsigned short>*/,
+                0/*addWeighted_gpu<short, unsigned char, short >*/,
+                0/*addWeighted_gpu<short, unsigned char, int   >*/,
+                0/*addWeighted_gpu<short, unsigned char, float >*/,
+                0/*addWeighted_gpu<short, unsigned char, double>*/
+            },
+            {
+                0/*addWeighted_gpu<short, signed char, unsigned char >*/,
+                0/*addWeighted_gpu<short, signed char, signed char >*/,
+                0/*addWeighted_gpu<short, signed char, unsigned short>*/,
+                0/*addWeighted_gpu<short, signed char, short >*/,
+                0/*addWeighted_gpu<short, signed char, int   >*/,
+                0/*addWeighted_gpu<short, signed char, float >*/,
+                0/*addWeighted_gpu<short, signed char, double>*/
+            },
+            {
+                0/*addWeighted_gpu<short, unsigned short, unsigned char >*/,
+                0/*addWeighted_gpu<short, unsigned short, signed char >*/,
+                0/*addWeighted_gpu<short, unsigned short, unsigned short>*/,
+                0/*addWeighted_gpu<short, unsigned short, short >*/,
+                0/*addWeighted_gpu<short, unsigned short, int   >*/,
+                0/*addWeighted_gpu<short, unsigned short, float >*/,
+                0/*addWeighted_gpu<short, unsigned short, double>*/
+            },
+            {
+                addWeighted_gpu<short, short, unsigned char >,
+                addWeighted_gpu<short, short, signed char >,
+                addWeighted_gpu<short, short, unsigned short>,
+                addWeighted_gpu<short, short, short >,
+                addWeighted_gpu<short, short, int   >,
+                addWeighted_gpu<short, short, float >,
+                addWeighted_gpu<short, short, double>
+            },
+            {
+                addWeighted_gpu<short, int, unsigned char >,
+                addWeighted_gpu<short, int, signed char >,
+                addWeighted_gpu<short, int, unsigned short>,
+                addWeighted_gpu<short, int, short >,
+                addWeighted_gpu<short, int, int   >,
+                addWeighted_gpu<short, int, float >,
+                addWeighted_gpu<short, int, double>
+            },
+            {
+                addWeighted_gpu<short, float, unsigned char >,
+                addWeighted_gpu<short, float, signed char >,
+                addWeighted_gpu<short, float, unsigned short>,
+                addWeighted_gpu<short, float, short >,
+                addWeighted_gpu<short, float, int   >,
+                addWeighted_gpu<short, float, float >,
+                addWeighted_gpu<short, float, double>
+            },
+            {
+                addWeighted_gpu<short, double, unsigned char >,
+                addWeighted_gpu<short, double, signed char >,
+                addWeighted_gpu<short, double, unsigned short>,
+                addWeighted_gpu<short, double, short >,
+                addWeighted_gpu<short, double, int   >,
+                addWeighted_gpu<short, double, float >,
+                addWeighted_gpu<short, double, double>
+            }
+        },
+        {
+            {
+                0/*addWeighted_gpu<int, unsigned char, unsigned char >*/,
+                0/*addWeighted_gpu<int, unsigned char, signed char >*/,
+                0/*addWeighted_gpu<int, unsigned char, unsigned short>*/,
+                0/*addWeighted_gpu<int, unsigned char, short >*/,
+                0/*addWeighted_gpu<int, unsigned char, int   >*/,
+                0/*addWeighted_gpu<int, unsigned char, float >*/,
+                0/*addWeighted_gpu<int, unsigned char, double>*/
+            },
+            {
+                0/*addWeighted_gpu<int, signed char, unsigned char >*/,
+                0/*addWeighted_gpu<int, signed char, signed char >*/,
+                0/*addWeighted_gpu<int, signed char, unsigned short>*/,
+                0/*addWeighted_gpu<int, signed char, short >*/,
+                0/*addWeighted_gpu<int, signed char, int   >*/,
+                0/*addWeighted_gpu<int, signed char, float >*/,
+                0/*addWeighted_gpu<int, signed char, double>*/
+            },
+            {
+                0/*addWeighted_gpu<int, unsigned short, unsigned char >*/,
+                0/*addWeighted_gpu<int, unsigned short, signed char >*/,
+                0/*addWeighted_gpu<int, unsigned short, unsigned short>*/,
+                0/*addWeighted_gpu<int, unsigned short, short >*/,
+                0/*addWeighted_gpu<int, unsigned short, int   >*/,
+                0/*addWeighted_gpu<int, unsigned short, float >*/,
+                0/*addWeighted_gpu<int, unsigned short, double>*/
+            },
+            {
+                0/*addWeighted_gpu<int, short, unsigned char >*/,
+                0/*addWeighted_gpu<int, short, signed char >*/,
+                0/*addWeighted_gpu<int, short, unsigned short>*/,
+                0/*addWeighted_gpu<int, short, short >*/,
+                0/*addWeighted_gpu<int, short, int   >*/,
+                0/*addWeighted_gpu<int, short, float >*/,
+                0/*addWeighted_gpu<int, short, double>*/
+            },
+            {
+                addWeighted_gpu<int, int, unsigned char >,
+                addWeighted_gpu<int, int, signed char >,
+                addWeighted_gpu<int, int, unsigned short>,
+                addWeighted_gpu<int, int, short >,
+                addWeighted_gpu<int, int, int   >,
+                addWeighted_gpu<int, int, float >,
+                addWeighted_gpu<int, int, double>
+            },
+            {
+                addWeighted_gpu<int, float, unsigned char >,
+                addWeighted_gpu<int, float, signed char >,
+                addWeighted_gpu<int, float, unsigned short>,
+                addWeighted_gpu<int, float, short >,
+                addWeighted_gpu<int, float, int   >,
+                addWeighted_gpu<int, float, float >,
+                addWeighted_gpu<int, float, double>
+            },
+            {
+                addWeighted_gpu<int, double, unsigned char >,
+                addWeighted_gpu<int, double, signed char >,
+                addWeighted_gpu<int, double, unsigned short>,
+                addWeighted_gpu<int, double, short >,
+                addWeighted_gpu<int, double, int   >,
+                addWeighted_gpu<int, double, float >,
+                addWeighted_gpu<int, double, double>
+            }
+        },
+        {
+            {
+                0/*addWeighted_gpu<float, unsigned char, unsigned char >*/,
+                0/*addWeighted_gpu<float, unsigned char, signed char >*/,
+                0/*addWeighted_gpu<float, unsigned char, unsigned short>*/,
+                0/*addWeighted_gpu<float, unsigned char, short >*/,
+                0/*addWeighted_gpu<float, unsigned char, int   >*/,
+                0/*addWeighted_gpu<float, unsigned char, float >*/,
+                0/*addWeighted_gpu<float, unsigned char, double>*/
+            },
+            {
+                0/*addWeighted_gpu<float, signed char, unsigned char >*/,
+                0/*addWeighted_gpu<float, signed char, signed char >*/,
+                0/*addWeighted_gpu<float, signed char, unsigned short>*/,
+                0/*addWeighted_gpu<float, signed char, short >*/,
+                0/*addWeighted_gpu<float, signed char, int   >*/,
+                0/*addWeighted_gpu<float, signed char, float >*/,
+                0/*addWeighted_gpu<float, signed char, double>*/
+            },
+            {
+                0/*addWeighted_gpu<float, unsigned short, unsigned char >*/,
+                0/*addWeighted_gpu<float, unsigned short, signed char >*/,
+                0/*addWeighted_gpu<float, unsigned short, unsigned short>*/,
+                0/*addWeighted_gpu<float, unsigned short, short >*/,
+                0/*addWeighted_gpu<float, unsigned short, int   >*/,
+                0/*addWeighted_gpu<float, unsigned short, float >*/,
+                0/*addWeighted_gpu<float, unsigned short, double>*/
+            },
+            {
+                0/*addWeighted_gpu<float, short, unsigned char >*/,
+                0/*addWeighted_gpu<float, short, signed char >*/,
+                0/*addWeighted_gpu<float, short, unsigned short>*/,
+                0/*addWeighted_gpu<float, short, short >*/,
+                0/*addWeighted_gpu<float, short, int   >*/,
+                0/*addWeighted_gpu<float, short, float >*/,
+                0/*addWeighted_gpu<float, short, double>*/
+            },
+            {
+                0/*addWeighted_gpu<float, int, unsigned char >*/,
+                0/*addWeighted_gpu<float, int, signed char >*/,
+                0/*addWeighted_gpu<float, int, unsigned short>*/,
+                0/*addWeighted_gpu<float, int, short >*/,
+                0/*addWeighted_gpu<float, int, int   >*/,
+                0/*addWeighted_gpu<float, int, float >*/,
+                0/*addWeighted_gpu<float, int, double>*/
+            },
+            {
+                addWeighted_gpu<float, float, unsigned char >,
+                addWeighted_gpu<float, float, signed char >,
+                addWeighted_gpu<float, float, unsigned short>,
+                addWeighted_gpu<float, float, short >,
+                addWeighted_gpu<float, float, int   >,
+                addWeighted_gpu<float, float, float >,
+                addWeighted_gpu<float, float, double>
+            },
+            {
+                addWeighted_gpu<float, double, unsigned char >,
+                addWeighted_gpu<float, double, signed char >,
+                addWeighted_gpu<float, double, unsigned short>,
+                addWeighted_gpu<float, double, short >,
+                addWeighted_gpu<float, double, int   >,
+                addWeighted_gpu<float, double, float >,
+                addWeighted_gpu<float, double, double>
+            }
+        },
+        {
+            {
+                0/*addWeighted_gpu<double, unsigned char, unsigned char >*/,
+                0/*addWeighted_gpu<double, unsigned char, signed char >*/,
+                0/*addWeighted_gpu<double, unsigned char, unsigned short>*/,
+                0/*addWeighted_gpu<double, unsigned char, short >*/,
+                0/*addWeighted_gpu<double, unsigned char, int   >*/,
+                0/*addWeighted_gpu<double, unsigned char, float >*/,
+                0/*addWeighted_gpu<double, unsigned char, double>*/
+            },
+            {
+                0/*addWeighted_gpu<double, signed char, unsigned char >*/,
+                0/*addWeighted_gpu<double, signed char, signed char >*/,
+                0/*addWeighted_gpu<double, signed char, unsigned short>*/,
+                0/*addWeighted_gpu<double, signed char, short >*/,
+                0/*addWeighted_gpu<double, signed char, int   >*/,
+                0/*addWeighted_gpu<double, signed char, float >*/,
+                0/*addWeighted_gpu<double, signed char, double>*/
+            },
+            {
+                0/*addWeighted_gpu<double, unsigned short, unsigned char >*/,
+                0/*addWeighted_gpu<double, unsigned short, signed char >*/,
+                0/*addWeighted_gpu<double, unsigned short, unsigned short>*/,
+                0/*addWeighted_gpu<double, unsigned short, short >*/,
+                0/*addWeighted_gpu<double, unsigned short, int   >*/,
+                0/*addWeighted_gpu<double, unsigned short, float >*/,
+                0/*addWeighted_gpu<double, unsigned short, double>*/
+            },
+            {
+                0/*addWeighted_gpu<double, short, unsigned char >*/,
+                0/*addWeighted_gpu<double, short, signed char >*/,
+                0/*addWeighted_gpu<double, short, unsigned short>*/,
+                0/*addWeighted_gpu<double, short, short >*/,
+                0/*addWeighted_gpu<double, short, int   >*/,
+                0/*addWeighted_gpu<double, short, float >*/,
+                0/*addWeighted_gpu<double, short, double>*/
+            },
+            {
+                0/*addWeighted_gpu<double, int, unsigned char >*/,
+                0/*addWeighted_gpu<double, int, signed char >*/,
+                0/*addWeighted_gpu<double, int, unsigned short>*/,
+                0/*addWeighted_gpu<double, int, short >*/,
+                0/*addWeighted_gpu<double, int, int   >*/,
+                0/*addWeighted_gpu<double, int, float >*/,
+                0/*addWeighted_gpu<double, int, double>*/
+            },
+            {
+                0/*addWeighted_gpu<double, float, unsigned char >*/,
+                0/*addWeighted_gpu<double, float, signed char >*/,
+                0/*addWeighted_gpu<double, float, unsigned short>*/,
+                0/*addWeighted_gpu<double, float, short >*/,
+                0/*addWeighted_gpu<double, float, int   >*/,
+                0/*addWeighted_gpu<double, float, float >*/,
+                0/*addWeighted_gpu<double, float, double>*/
+            },
+            {
+                addWeighted_gpu<double, double, unsigned char >,
+                addWeighted_gpu<double, double, signed char >,
+                addWeighted_gpu<double, double, unsigned short>,
+                addWeighted_gpu<double, double, short >,
+                addWeighted_gpu<double, double, int   >,
+                addWeighted_gpu<double, double, float >,
+                addWeighted_gpu<double, double, double>
+            }
+        }
+    };
+
+    callers[psrc1->depth()][psrc2->depth()][dst.depth()](psrc1->reshape(1), alpha, psrc2->reshape(1), beta, gamma, dst.reshape(1), StreamAccessor::getStream(stream));
+}
+
 #endif
diff --git a/modules/gpu/test/test_arithm.cpp b/modules/gpu/test/test_arithm.cpp
index a8fcccdca..31365ca1f 100644
--- a/modules/gpu/test/test_arithm.cpp
+++ b/modules/gpu/test/test_arithm.cpp
@@ -1135,7 +1135,7 @@ TEST_P(MinMax, Accuracy)
         return;
 
     PRINT_PARAM(devInfo);
-    PRINT_TYPE(type)
+    PRINT_TYPE(type);
     PRINT_PARAM(size);
 
     double minVal, maxVal;
@@ -1216,7 +1216,7 @@ TEST_P(MinMaxLoc, Accuracy)
         return;
 
     PRINT_PARAM(devInfo);
-    PRINT_TYPE(type)
+    PRINT_TYPE(type);
     PRINT_PARAM(size);
 
     double minVal, maxVal;
@@ -1281,7 +1281,7 @@ TEST_P(CountNonZero, Accuracy)
         return;
 
     PRINT_PARAM(devInfo);
-    PRINT_TYPE(type)
+    PRINT_TYPE(type);
     PRINT_PARAM(size);
 
     int n;
@@ -1333,7 +1333,7 @@ TEST_P(Sum, Accuracy)
         return;
 
     PRINT_PARAM(devInfo);
-    PRINT_TYPE(type)
+    PRINT_TYPE(type);
     PRINT_PARAM(size);
 
     cv::Scalar sum;
@@ -1385,7 +1385,7 @@ TEST_P(AbsSum, Accuracy)
         return;
 
     PRINT_PARAM(devInfo);
-    PRINT_TYPE(type)
+    PRINT_TYPE(type);
     PRINT_PARAM(size);
 
     cv::Scalar sum;
@@ -1439,7 +1439,7 @@ TEST_P(SqrSum, Accuracy)
         return;
 
     PRINT_PARAM(devInfo);
-    PRINT_TYPE(type)
+    PRINT_TYPE(type);
     PRINT_PARAM(size);
 
     cv::Scalar sum;
@@ -1500,7 +1500,7 @@ TEST_P(BitwiseNot, Accuracy)
         return;
 
     PRINT_PARAM(devInfo);
-    PRINT_TYPE(type)
+    PRINT_TYPE(type);
     PRINT_PARAM(size);
 
     cv::Mat dst;
@@ -1564,7 +1564,7 @@ TEST_P(BitwiseOr, Accuracy)
         return;
 
     PRINT_PARAM(devInfo);
-    PRINT_TYPE(type)
+    PRINT_TYPE(type);
     PRINT_PARAM(size);
 
     cv::Mat dst;
@@ -1628,7 +1628,7 @@ TEST_P(BitwiseAnd, Accuracy)
         return;
 
     PRINT_PARAM(devInfo);
-    PRINT_TYPE(type)
+    PRINT_TYPE(type);
     PRINT_PARAM(size);
 
     cv::Mat dst;
@@ -1692,7 +1692,7 @@ TEST_P(BitwiseXor, Accuracy)
         return;
 
     PRINT_PARAM(devInfo);
-    PRINT_TYPE(type)
+    PRINT_TYPE(type);
     PRINT_PARAM(size);
 
     cv::Mat dst;
@@ -1712,4 +1712,80 @@ INSTANTIATE_TEST_CASE_P(Arithm, BitwiseXor, testing::Combine(
                         testing::ValuesIn(devices()),
                         testing::ValuesIn(all_types())));
 
+//////////////////////////////////////////////////////////////////////////////
+// addWeighted
+
+struct AddWeighted : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int, int, int> >
+{
+    cv::gpu::DeviceInfo devInfo;
+    int type1;
+    int type2;
+    int dtype;
+
+    cv::Size size;
+    cv::Mat src1;
+    cv::Mat src2;
+    double alpha;
+    double beta;
+    double gamma;
+
+    cv::Mat dst_gold;
+
+    virtual void SetUp() 
+    {
+        devInfo = std::tr1::get<0>(GetParam());
+        type1 = std::tr1::get<1>(GetParam());
+        type2 = std::tr1::get<2>(GetParam());
+        dtype = std::tr1::get<3>(GetParam());
+
+        cv::gpu::setDevice(devInfo.deviceID());
+
+        cv::RNG& rng = cvtest::TS::ptr()->get_rng();
+
+        size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200));
+
+        src1 = cvtest::randomMat(rng, size, type1, 0.0, 255.0, false);
+        src2 = cvtest::randomMat(rng, size, type2, 0.0, 255.0, false);
+
+        alpha = rng.uniform(-10.0, 10.0);
+        beta = rng.uniform(-10.0, 10.0);
+        gamma = rng.uniform(-10.0, 10.0);
+
+        cv::addWeighted(src1, alpha, src2, beta, gamma, dst_gold, dtype);
+    }
+};
+
+TEST_P(AddWeighted, Accuracy) 
+{
+    if ((src1.depth() == CV_64F || src2.depth() == CV_64F || dst_gold.depth() == CV_64F) && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))
+        return;
+
+    PRINT_PARAM(devInfo);
+    PRINT_TYPE(type1);
+    PRINT_TYPE(type2);
+    PRINT_TYPE(dtype);
+    PRINT_PARAM(size);
+    PRINT_PARAM(alpha);
+    PRINT_PARAM(beta);
+    PRINT_PARAM(gamma);
+
+    cv::Mat dst;
+    
+    ASSERT_NO_THROW(
+        cv::gpu::GpuMat dev_dst;
+
+        cv::gpu::addWeighted(cv::gpu::GpuMat(src1), alpha, cv::gpu::GpuMat(src2), beta, gamma, dev_dst, dtype);
+
+        dev_dst.download(dst);
+    );
+
+    EXPECT_MAT_NEAR(dst_gold, dst, dtype < CV_32F ? 1.0 : 1e-12);
+}
+
+INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, testing::Combine(
+                        testing::ValuesIn(devices()),
+                        testing::ValuesIn(types(CV_8U, CV_64F, 1, 1)),
+                        testing::ValuesIn(types(CV_8U, CV_64F, 1, 1)),
+                        testing::ValuesIn(types(CV_8U, CV_64F, 1, 1))));
+
 #endif // HAVE_CUDA