implemented gpu::addWeighted

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index dbb1107..7377f52 100644 (file)
--- 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))\r
         CV_EXPORTS void max(const GpuMat& src1, double src2, GpuMat& dst, Stream& stream = Stream::Null());\r
 \r
+        //! computes the weighted sum of two arrays\r
+        CV_EXPORTS void addWeighted(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst, \r
+            int dtype = -1, Stream& stream = Stream::Null());\r
+\r
 \r
         ////////////////////////////// Image processing //////////////////////////////\r
 \r

diff --git a/modules/gpu/perf/perf_arithm.cpp b/modules/gpu/perf/perf_arithm.cpp
index f5c4da3..2c91590 100644 (file)
--- 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
 \r
     SANITY_CHECK(dst);\r
 }\r
+\r
+PERF_TEST_P(DevInfo_Size_MatType, addWeighted, testing::Combine(testing::ValuesIn(devices()), \r
+                                                                testing::Values(GPU_TYPICAL_MAT_SIZES), \r
+                                                                testing::Values(CV_8UC1, CV_16UC1, CV_32FC1)))\r
+{\r
+    DeviceInfo devInfo = std::tr1::get<0>(GetParam());\r
+    Size size = std::tr1::get<1>(GetParam());\r
+    int type = std::tr1::get<2>(GetParam());\r
+\r
+    setDevice(devInfo.deviceID());\r
+\r
+    Mat src1_host(size, type);\r
+    Mat src2_host(size, type);\r
+\r
+    declare.in(src1_host, src2_host, WARMUP_RNG);\r
+\r
+    GpuMat src1(src1_host);\r
+    GpuMat src2(src2_host);\r
+    GpuMat dst(size, type);\r
+\r
+    declare.time(0.5).iterations(100);\r
+\r
+    SIMPLE_TEST_CYCLE()\r
+    {\r
+        addWeighted(src1, 0.5, src2, 0.5, 0.0, dst);\r
+    }\r
+\r
+    Mat dst_host = dst;\r
+\r
+    SANITY_CHECK(dst_host);\r
+}\r

diff --git a/modules/gpu/src/cuda/element_operations.cu b/modules/gpu/src/cuda/element_operations.cu
index 71574b7..7548e90 100644 (file)
--- 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
     }\r
 \r
     template void multiplyScalar_gpu<uchar, uchar>(const DevMem2D& src, float scale, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+\r
+    //////////////////////////////////////////////////////////////////////////\r
+    // addWeighted\r
+\r
+    template <typename T1, typename T2, typename D> struct AddWeighted : binary_function<T1, T2, D>\r
+    {\r
+        __host__ __device__ __forceinline__ AddWeighted(double alpha_, double beta_, double gamma_) : alpha(alpha_), beta(beta_), gamma(gamma_) {}\r
+\r
+        __device__ __forceinline__ D operator ()(typename TypeTraits<T1>::ParameterType a, typename TypeTraits<T2>::ParameterType b) const\r
+        {\r
+            return saturate_cast<D>(alpha * a + beta * b + gamma);\r
+        }\r
+\r
+        const double alpha;\r
+        const double beta;\r
+        const double gamma;\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< AddWeighted<ushort, ushort, ushort> > : DefaultTransformFunctorTraits< AddWeighted<ushort, ushort, ushort> >\r
+    {\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddWeighted<ushort, ushort, short> > : DefaultTransformFunctorTraits< AddWeighted<ushort, ushort, short> >\r
+    {\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddWeighted<ushort, short, ushort> > : DefaultTransformFunctorTraits< AddWeighted<ushort, short, ushort> >\r
+    {\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddWeighted<ushort, short, short> > : DefaultTransformFunctorTraits< AddWeighted<ushort, short, short> >\r
+    {\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddWeighted<short, short, ushort> > : DefaultTransformFunctorTraits< AddWeighted<short, short, ushort> >\r
+    {\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddWeighted<short, short, short> > : DefaultTransformFunctorTraits< AddWeighted<short, short, short> >\r
+    {\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< AddWeighted<int, int, int> > : DefaultTransformFunctorTraits< AddWeighted<int, int, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddWeighted<int, int, float> > : DefaultTransformFunctorTraits< AddWeighted<int, int, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddWeighted<int, float, int> > : DefaultTransformFunctorTraits< AddWeighted<int, float, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddWeighted<int, float, float> > : DefaultTransformFunctorTraits< AddWeighted<int, float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddWeighted<float, float, int> > : DefaultTransformFunctorTraits< AddWeighted<float, float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddWeighted<float, float, float> > : DefaultTransformFunctorTraits< AddWeighted<float, float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T1, typename T2, typename D>\r
+    void addWeighted_gpu(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall( cudaSetDoubleForDevice(&alpha) );\r
+        cudaSafeCall( cudaSetDoubleForDevice(&beta) );\r
+        cudaSafeCall( cudaSetDoubleForDevice(&gamma) );\r
+\r
+        AddWeighted<T1, T2, D> op(alpha, beta, gamma);\r
+\r
+        transform(static_cast< DevMem2D_<T1> >(src1), static_cast< DevMem2D_<T2> >(src2), static_cast< DevMem2D_<D> >(dst), op, stream);\r
+    }\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    \r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    \r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+\r
+    \r
+\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
+    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);\r
 }}}\r

diff --git a/modules/gpu/src/element_operations.cpp b/modules/gpu/src/element_operations.cpp
index 68724fe..4cf857e 100644 (file)
--- 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(); }\r
 void cv::gpu::max(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }\r
 double cv::gpu::threshold(const GpuMat&, GpuMat&, double, double, int, Stream&) {throw_nogpu(); return 0.0;}\r
-void cv::gpu::pow(const GpuMat&, double, GpuMat&, Stream&)  { throw_nogpu(); }\r
+void cv::gpu::pow(const GpuMat&, double, GpuMat&, Stream&) { throw_nogpu(); }\r
+void cv::gpu::addWeighted(const GpuMat&, double, const GpuMat&, double, double, GpuMat&, int, Stream&) { throw_nogpu(); }\r
 \r
 #else\r
 \r
@@ -248,7 +249,7 @@ void cv::gpu::multiply(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream&
     {\r
         dst.create(src.size(), src.type());\r
 \r
-        device::multiplyScalar_gpu<uchar, uchar>(src.reshape(1), (float)(sc[0]), dst, StreamAccessor::getStream(stream));\r
+        device::multiplyScalar_gpu<unsigned char, unsigned char>(src.reshape(1), (float)(sc[0]), dst, StreamAccessor::getStream(stream));\r
     }\r
     else\r
     {\r
@@ -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);\r
     static const func_t funcs[] = \r
     {\r
-        min_caller<uchar>, min_caller<schar>, min_caller<ushort>, min_caller<short>, min_caller<int>, \r
+        min_caller<unsigned char>, min_caller<signed char>, min_caller<unsigned short>, min_caller<short>, min_caller<int>, \r
         min_caller<float>, min_caller<double>\r
     };\r
     funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));\r
@@ -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);\r
     static const func_t funcs[] = \r
     {\r
-        min_caller<uchar>, min_caller<schar>, min_caller<ushort>, min_caller<short>, min_caller<int>, \r
+        min_caller<unsigned char>, min_caller<signed char>, min_caller<unsigned short>, min_caller<short>, min_caller<int>, \r
         min_caller<float>, min_caller<double>\r
     };\r
     funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));\r
@@ -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);\r
     static const func_t funcs[] = \r
     {\r
-        max_caller<uchar>, max_caller<schar>, max_caller<ushort>, max_caller<short>, max_caller<int>, \r
+        max_caller<unsigned char>, max_caller<signed char>, max_caller<unsigned short>, max_caller<short>, max_caller<int>, \r
         max_caller<float>, max_caller<double>\r
     };\r
     funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));\r
@@ -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);\r
     static const func_t funcs[] = \r
     {\r
-        max_caller<uchar>, max_caller<schar>, max_caller<ushort>, max_caller<short>, max_caller<int>, \r
+        max_caller<unsigned char>, max_caller<signed char>, max_caller<unsigned short>, max_caller<short>, max_caller<int>, \r
         max_caller<float>, max_caller<double>\r
     };\r
     funcs[src1.depth()](src1, src2, dst, StreamAccessor::getStream(stream));\r
@@ -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));    \r
 }\r
 \r
+////////////////////////////////////////////////////////////////////////\r
+// addWeighted\r
+\r
+namespace cv { namespace gpu { namespace device\r
+{\r
+    template <typename T1, typename T2, typename D>\r
+    void addWeighted_gpu(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);\r
+}}}\r
+\r
+void cv::gpu::addWeighted(const GpuMat& src1, double alpha, const GpuMat& src2, double beta, double gamma, GpuMat& dst, int dtype, Stream& stream)\r
+{\r
+    CV_Assert(src1.size() == src2.size());\r
+    CV_Assert(src1.type() == src2.type() || (dtype >= 0 && src1.channels() == src2.channels()));\r
+\r
+    dtype = dtype >= 0 ? CV_MAKETYPE(dtype, src1.channels()) : src1.type();\r
+\r
+    dst.create(src1.size(), dtype);\r
+\r
+    const GpuMat* psrc1 = &src1;\r
+    const GpuMat* psrc2 = &src2;\r
+\r
+    if (src1.depth() > src2.depth())\r
+    {\r
+        std::swap(psrc1, psrc2);\r
+        std::swap(alpha, beta);\r
+    }\r
+\r
+    typedef void (*caller_t)(const DevMem2D& src1, double alpha, const DevMem2D& src2, double beta, double gamma, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+    using namespace cv::gpu::device;\r
+\r
+    static const caller_t callers[7][7][7] =\r
+    {\r
+        {\r
+            {\r
+                addWeighted_gpu<unsigned char, unsigned char, unsigned char >,\r
+                addWeighted_gpu<unsigned char, unsigned char, signed char >,\r
+                addWeighted_gpu<unsigned char, unsigned char, unsigned short>,\r
+                addWeighted_gpu<unsigned char, unsigned char, short >,\r
+                addWeighted_gpu<unsigned char, unsigned char, int   >,\r
+                addWeighted_gpu<unsigned char, unsigned char, float >,\r
+                addWeighted_gpu<unsigned char, unsigned char, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned char, signed char, unsigned char >,\r
+                addWeighted_gpu<unsigned char, signed char, signed char >,\r
+                addWeighted_gpu<unsigned char, signed char, unsigned short>,\r
+                addWeighted_gpu<unsigned char, signed char, short >,\r
+                addWeighted_gpu<unsigned char, signed char, int   >,\r
+                addWeighted_gpu<unsigned char, signed char, float >,\r
+                addWeighted_gpu<unsigned char, signed char, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned char, unsigned short, unsigned char >,\r
+                addWeighted_gpu<unsigned char, unsigned short, signed char >,\r
+                addWeighted_gpu<unsigned char, unsigned short, unsigned short>,\r
+                addWeighted_gpu<unsigned char, unsigned short, short >,\r
+                addWeighted_gpu<unsigned char, unsigned short, int   >,\r
+                addWeighted_gpu<unsigned char, unsigned short, float >,\r
+                addWeighted_gpu<unsigned char, unsigned short, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned char, short, unsigned char >,\r
+                addWeighted_gpu<unsigned char, short, signed char >,\r
+                addWeighted_gpu<unsigned char, short, unsigned short>,\r
+                addWeighted_gpu<unsigned char, short, short >,\r
+                addWeighted_gpu<unsigned char, short, int   >,\r
+                addWeighted_gpu<unsigned char, short, float >,\r
+                addWeighted_gpu<unsigned char, short, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned char, int, unsigned char >,\r
+                addWeighted_gpu<unsigned char, int, signed char >,\r
+                addWeighted_gpu<unsigned char, int, unsigned short>,\r
+                addWeighted_gpu<unsigned char, int, short >,\r
+                addWeighted_gpu<unsigned char, int, int   >,\r
+                addWeighted_gpu<unsigned char, int, float >,\r
+                addWeighted_gpu<unsigned char, int, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned char, float, unsigned char >,\r
+                addWeighted_gpu<unsigned char, float, signed char >,\r
+                addWeighted_gpu<unsigned char, float, unsigned short>,\r
+                addWeighted_gpu<unsigned char, float, short >,\r
+                addWeighted_gpu<unsigned char, float, int   >,\r
+                addWeighted_gpu<unsigned char, float, float >,\r
+                addWeighted_gpu<unsigned char, float, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned char, double, unsigned char >,\r
+                addWeighted_gpu<unsigned char, double, signed char >,\r
+                addWeighted_gpu<unsigned char, double, unsigned short>,\r
+                addWeighted_gpu<unsigned char, double, short >,\r
+                addWeighted_gpu<unsigned char, double, int   >,\r
+                addWeighted_gpu<unsigned char, double, float >,\r
+                addWeighted_gpu<unsigned char, double, double>\r
+            }\r
+        },\r
+        {\r
+            {\r
+                0/*addWeighted_gpu<signed char, unsigned char, unsigned char >*/,\r
+                0/*addWeighted_gpu<signed char, unsigned char, signed char >*/,\r
+                0/*addWeighted_gpu<signed char, unsigned char, unsigned short>*/,\r
+                0/*addWeighted_gpu<signed char, unsigned char, short >*/,\r
+                0/*addWeighted_gpu<signed char, unsigned char, int   >*/,\r
+                0/*addWeighted_gpu<signed char, unsigned char, float >*/,\r
+                0/*addWeighted_gpu<signed char, unsigned char, double>*/\r
+            },\r
+            {\r
+                addWeighted_gpu<signed char, signed char, unsigned char >,\r
+                addWeighted_gpu<signed char, signed char, signed char >,\r
+                addWeighted_gpu<signed char, signed char, unsigned short>,\r
+                addWeighted_gpu<signed char, signed char, short >,\r
+                addWeighted_gpu<signed char, signed char, int   >,\r
+                addWeighted_gpu<signed char, signed char, float >,\r
+                addWeighted_gpu<signed char, signed char, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<signed char, unsigned short, unsigned char >,\r
+                addWeighted_gpu<signed char, unsigned short, signed char >,\r
+                addWeighted_gpu<signed char, unsigned short, unsigned short>,\r
+                addWeighted_gpu<signed char, unsigned short, short >,\r
+                addWeighted_gpu<signed char, unsigned short, int   >,\r
+                addWeighted_gpu<signed char, unsigned short, float >,\r
+                addWeighted_gpu<signed char, unsigned short, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<signed char, short, unsigned char >,\r
+                addWeighted_gpu<signed char, short, signed char >,\r
+                addWeighted_gpu<signed char, short, unsigned short>,\r
+                addWeighted_gpu<signed char, short, short >,\r
+                addWeighted_gpu<signed char, short, int   >,\r
+                addWeighted_gpu<signed char, short, float >,\r
+                addWeighted_gpu<signed char, short, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<signed char, int, unsigned char >,\r
+                addWeighted_gpu<signed char, int, signed char >,\r
+                addWeighted_gpu<signed char, int, unsigned short>,\r
+                addWeighted_gpu<signed char, int, short >,\r
+                addWeighted_gpu<signed char, int, int   >,\r
+                addWeighted_gpu<signed char, int, float >,\r
+                addWeighted_gpu<signed char, int, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<signed char, float, unsigned char >,\r
+                addWeighted_gpu<signed char, float, signed char >,\r
+                addWeighted_gpu<signed char, float, unsigned short>,\r
+                addWeighted_gpu<signed char, float, short >,\r
+                addWeighted_gpu<signed char, float, int   >,\r
+                addWeighted_gpu<signed char, float, float >,\r
+                addWeighted_gpu<signed char, float, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<signed char, double, unsigned char >,\r
+                addWeighted_gpu<signed char, double, signed char >,\r
+                addWeighted_gpu<signed char, double, unsigned short>,\r
+                addWeighted_gpu<signed char, double, short >,\r
+                addWeighted_gpu<signed char, double, int   >,\r
+                addWeighted_gpu<signed char, double, float >,\r
+                addWeighted_gpu<signed char, double, double>\r
+            }\r
+        },\r
+        {\r
+            {\r
+                0/*addWeighted_gpu<unsigned short, unsigned char, unsigned char >*/,\r
+                0/*addWeighted_gpu<unsigned short, unsigned char, signed char >*/,\r
+                0/*addWeighted_gpu<unsigned short, unsigned char, unsigned short>*/,\r
+                0/*addWeighted_gpu<unsigned short, unsigned char, short >*/,\r
+                0/*addWeighted_gpu<unsigned short, unsigned char, int   >*/,\r
+                0/*addWeighted_gpu<unsigned short, unsigned char, float >*/,\r
+                0/*addWeighted_gpu<unsigned short, unsigned char, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<unsigned short, signed char, unsigned char >*/,\r
+                0/*addWeighted_gpu<unsigned short, signed char, signed char >*/,\r
+                0/*addWeighted_gpu<unsigned short, signed char, unsigned short>*/,\r
+                0/*addWeighted_gpu<unsigned short, signed char, short >*/,\r
+                0/*addWeighted_gpu<unsigned short, signed char, int   >*/,\r
+                0/*addWeighted_gpu<unsigned short, signed char, float >*/,\r
+                0/*addWeighted_gpu<unsigned short, signed char, double>*/\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned short, unsigned short, unsigned char >,\r
+                addWeighted_gpu<unsigned short, unsigned short, signed char >,\r
+                addWeighted_gpu<unsigned short, unsigned short, unsigned short>,\r
+                addWeighted_gpu<unsigned short, unsigned short, short >,\r
+                addWeighted_gpu<unsigned short, unsigned short, int   >,\r
+                addWeighted_gpu<unsigned short, unsigned short, float >,\r
+                addWeighted_gpu<unsigned short, unsigned short, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned short, short, unsigned char >,\r
+                addWeighted_gpu<unsigned short, short, signed char >,\r
+                addWeighted_gpu<unsigned short, short, unsigned short>,\r
+                addWeighted_gpu<unsigned short, short, short >,\r
+                addWeighted_gpu<unsigned short, short, int   >,\r
+                addWeighted_gpu<unsigned short, short, float >,\r
+                addWeighted_gpu<unsigned short, short, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned short, int, unsigned char >,\r
+                addWeighted_gpu<unsigned short, int, signed char >,\r
+                addWeighted_gpu<unsigned short, int, unsigned short>,\r
+                addWeighted_gpu<unsigned short, int, short >,\r
+                addWeighted_gpu<unsigned short, int, int   >,\r
+                addWeighted_gpu<unsigned short, int, float >,\r
+                addWeighted_gpu<unsigned short, int, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned short, float, unsigned char >,\r
+                addWeighted_gpu<unsigned short, float, signed char >,\r
+                addWeighted_gpu<unsigned short, float, unsigned short>,\r
+                addWeighted_gpu<unsigned short, float, short >,\r
+                addWeighted_gpu<unsigned short, float, int   >,\r
+                addWeighted_gpu<unsigned short, float, float >,\r
+                addWeighted_gpu<unsigned short, float, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<unsigned short, double, unsigned char >,\r
+                addWeighted_gpu<unsigned short, double, signed char >,\r
+                addWeighted_gpu<unsigned short, double, unsigned short>,\r
+                addWeighted_gpu<unsigned short, double, short >,\r
+                addWeighted_gpu<unsigned short, double, int   >,\r
+                addWeighted_gpu<unsigned short, double, float >,\r
+                addWeighted_gpu<unsigned short, double, double>\r
+            }\r
+        },\r
+        {\r
+            {\r
+                0/*addWeighted_gpu<short, unsigned char, unsigned char >*/,\r
+                0/*addWeighted_gpu<short, unsigned char, signed char >*/,\r
+                0/*addWeighted_gpu<short, unsigned char, unsigned short>*/,\r
+                0/*addWeighted_gpu<short, unsigned char, short >*/,\r
+                0/*addWeighted_gpu<short, unsigned char, int   >*/,\r
+                0/*addWeighted_gpu<short, unsigned char, float >*/,\r
+                0/*addWeighted_gpu<short, unsigned char, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<short, signed char, unsigned char >*/,\r
+                0/*addWeighted_gpu<short, signed char, signed char >*/,\r
+                0/*addWeighted_gpu<short, signed char, unsigned short>*/,\r
+                0/*addWeighted_gpu<short, signed char, short >*/,\r
+                0/*addWeighted_gpu<short, signed char, int   >*/,\r
+                0/*addWeighted_gpu<short, signed char, float >*/,\r
+                0/*addWeighted_gpu<short, signed char, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<short, unsigned short, unsigned char >*/,\r
+                0/*addWeighted_gpu<short, unsigned short, signed char >*/,\r
+                0/*addWeighted_gpu<short, unsigned short, unsigned short>*/,\r
+                0/*addWeighted_gpu<short, unsigned short, short >*/,\r
+                0/*addWeighted_gpu<short, unsigned short, int   >*/,\r
+                0/*addWeighted_gpu<short, unsigned short, float >*/,\r
+                0/*addWeighted_gpu<short, unsigned short, double>*/\r
+            },\r
+            {\r
+                addWeighted_gpu<short, short, unsigned char >,\r
+                addWeighted_gpu<short, short, signed char >,\r
+                addWeighted_gpu<short, short, unsigned short>,\r
+                addWeighted_gpu<short, short, short >,\r
+                addWeighted_gpu<short, short, int   >,\r
+                addWeighted_gpu<short, short, float >,\r
+                addWeighted_gpu<short, short, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<short, int, unsigned char >,\r
+                addWeighted_gpu<short, int, signed char >,\r
+                addWeighted_gpu<short, int, unsigned short>,\r
+                addWeighted_gpu<short, int, short >,\r
+                addWeighted_gpu<short, int, int   >,\r
+                addWeighted_gpu<short, int, float >,\r
+                addWeighted_gpu<short, int, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<short, float, unsigned char >,\r
+                addWeighted_gpu<short, float, signed char >,\r
+                addWeighted_gpu<short, float, unsigned short>,\r
+                addWeighted_gpu<short, float, short >,\r
+                addWeighted_gpu<short, float, int   >,\r
+                addWeighted_gpu<short, float, float >,\r
+                addWeighted_gpu<short, float, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<short, double, unsigned char >,\r
+                addWeighted_gpu<short, double, signed char >,\r
+                addWeighted_gpu<short, double, unsigned short>,\r
+                addWeighted_gpu<short, double, short >,\r
+                addWeighted_gpu<short, double, int   >,\r
+                addWeighted_gpu<short, double, float >,\r
+                addWeighted_gpu<short, double, double>\r
+            }\r
+        },\r
+        {\r
+            {\r
+                0/*addWeighted_gpu<int, unsigned char, unsigned char >*/,\r
+                0/*addWeighted_gpu<int, unsigned char, signed char >*/,\r
+                0/*addWeighted_gpu<int, unsigned char, unsigned short>*/,\r
+                0/*addWeighted_gpu<int, unsigned char, short >*/,\r
+                0/*addWeighted_gpu<int, unsigned char, int   >*/,\r
+                0/*addWeighted_gpu<int, unsigned char, float >*/,\r
+                0/*addWeighted_gpu<int, unsigned char, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<int, signed char, unsigned char >*/,\r
+                0/*addWeighted_gpu<int, signed char, signed char >*/,\r
+                0/*addWeighted_gpu<int, signed char, unsigned short>*/,\r
+                0/*addWeighted_gpu<int, signed char, short >*/,\r
+                0/*addWeighted_gpu<int, signed char, int   >*/,\r
+                0/*addWeighted_gpu<int, signed char, float >*/,\r
+                0/*addWeighted_gpu<int, signed char, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<int, unsigned short, unsigned char >*/,\r
+                0/*addWeighted_gpu<int, unsigned short, signed char >*/,\r
+                0/*addWeighted_gpu<int, unsigned short, unsigned short>*/,\r
+                0/*addWeighted_gpu<int, unsigned short, short >*/,\r
+                0/*addWeighted_gpu<int, unsigned short, int   >*/,\r
+                0/*addWeighted_gpu<int, unsigned short, float >*/,\r
+                0/*addWeighted_gpu<int, unsigned short, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<int, short, unsigned char >*/,\r
+                0/*addWeighted_gpu<int, short, signed char >*/,\r
+                0/*addWeighted_gpu<int, short, unsigned short>*/,\r
+                0/*addWeighted_gpu<int, short, short >*/,\r
+                0/*addWeighted_gpu<int, short, int   >*/,\r
+                0/*addWeighted_gpu<int, short, float >*/,\r
+                0/*addWeighted_gpu<int, short, double>*/\r
+            },\r
+            {\r
+                addWeighted_gpu<int, int, unsigned char >,\r
+                addWeighted_gpu<int, int, signed char >,\r
+                addWeighted_gpu<int, int, unsigned short>,\r
+                addWeighted_gpu<int, int, short >,\r
+                addWeighted_gpu<int, int, int   >,\r
+                addWeighted_gpu<int, int, float >,\r
+                addWeighted_gpu<int, int, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<int, float, unsigned char >,\r
+                addWeighted_gpu<int, float, signed char >,\r
+                addWeighted_gpu<int, float, unsigned short>,\r
+                addWeighted_gpu<int, float, short >,\r
+                addWeighted_gpu<int, float, int   >,\r
+                addWeighted_gpu<int, float, float >,\r
+                addWeighted_gpu<int, float, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<int, double, unsigned char >,\r
+                addWeighted_gpu<int, double, signed char >,\r
+                addWeighted_gpu<int, double, unsigned short>,\r
+                addWeighted_gpu<int, double, short >,\r
+                addWeighted_gpu<int, double, int   >,\r
+                addWeighted_gpu<int, double, float >,\r
+                addWeighted_gpu<int, double, double>\r
+            }\r
+        },\r
+        {\r
+            {\r
+                0/*addWeighted_gpu<float, unsigned char, unsigned char >*/,\r
+                0/*addWeighted_gpu<float, unsigned char, signed char >*/,\r
+                0/*addWeighted_gpu<float, unsigned char, unsigned short>*/,\r
+                0/*addWeighted_gpu<float, unsigned char, short >*/,\r
+                0/*addWeighted_gpu<float, unsigned char, int   >*/,\r
+                0/*addWeighted_gpu<float, unsigned char, float >*/,\r
+                0/*addWeighted_gpu<float, unsigned char, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<float, signed char, unsigned char >*/,\r
+                0/*addWeighted_gpu<float, signed char, signed char >*/,\r
+                0/*addWeighted_gpu<float, signed char, unsigned short>*/,\r
+                0/*addWeighted_gpu<float, signed char, short >*/,\r
+                0/*addWeighted_gpu<float, signed char, int   >*/,\r
+                0/*addWeighted_gpu<float, signed char, float >*/,\r
+                0/*addWeighted_gpu<float, signed char, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<float, unsigned short, unsigned char >*/,\r
+                0/*addWeighted_gpu<float, unsigned short, signed char >*/,\r
+                0/*addWeighted_gpu<float, unsigned short, unsigned short>*/,\r
+                0/*addWeighted_gpu<float, unsigned short, short >*/,\r
+                0/*addWeighted_gpu<float, unsigned short, int   >*/,\r
+                0/*addWeighted_gpu<float, unsigned short, float >*/,\r
+                0/*addWeighted_gpu<float, unsigned short, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<float, short, unsigned char >*/,\r
+                0/*addWeighted_gpu<float, short, signed char >*/,\r
+                0/*addWeighted_gpu<float, short, unsigned short>*/,\r
+                0/*addWeighted_gpu<float, short, short >*/,\r
+                0/*addWeighted_gpu<float, short, int   >*/,\r
+                0/*addWeighted_gpu<float, short, float >*/,\r
+                0/*addWeighted_gpu<float, short, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<float, int, unsigned char >*/,\r
+                0/*addWeighted_gpu<float, int, signed char >*/,\r
+                0/*addWeighted_gpu<float, int, unsigned short>*/,\r
+                0/*addWeighted_gpu<float, int, short >*/,\r
+                0/*addWeighted_gpu<float, int, int   >*/,\r
+                0/*addWeighted_gpu<float, int, float >*/,\r
+                0/*addWeighted_gpu<float, int, double>*/\r
+            },\r
+            {\r
+                addWeighted_gpu<float, float, unsigned char >,\r
+                addWeighted_gpu<float, float, signed char >,\r
+                addWeighted_gpu<float, float, unsigned short>,\r
+                addWeighted_gpu<float, float, short >,\r
+                addWeighted_gpu<float, float, int   >,\r
+                addWeighted_gpu<float, float, float >,\r
+                addWeighted_gpu<float, float, double>\r
+            },\r
+            {\r
+                addWeighted_gpu<float, double, unsigned char >,\r
+                addWeighted_gpu<float, double, signed char >,\r
+                addWeighted_gpu<float, double, unsigned short>,\r
+                addWeighted_gpu<float, double, short >,\r
+                addWeighted_gpu<float, double, int   >,\r
+                addWeighted_gpu<float, double, float >,\r
+                addWeighted_gpu<float, double, double>\r
+            }\r
+        },\r
+        {\r
+            {\r
+                0/*addWeighted_gpu<double, unsigned char, unsigned char >*/,\r
+                0/*addWeighted_gpu<double, unsigned char, signed char >*/,\r
+                0/*addWeighted_gpu<double, unsigned char, unsigned short>*/,\r
+                0/*addWeighted_gpu<double, unsigned char, short >*/,\r
+                0/*addWeighted_gpu<double, unsigned char, int   >*/,\r
+                0/*addWeighted_gpu<double, unsigned char, float >*/,\r
+                0/*addWeighted_gpu<double, unsigned char, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<double, signed char, unsigned char >*/,\r
+                0/*addWeighted_gpu<double, signed char, signed char >*/,\r
+                0/*addWeighted_gpu<double, signed char, unsigned short>*/,\r
+                0/*addWeighted_gpu<double, signed char, short >*/,\r
+                0/*addWeighted_gpu<double, signed char, int   >*/,\r
+                0/*addWeighted_gpu<double, signed char, float >*/,\r
+                0/*addWeighted_gpu<double, signed char, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<double, unsigned short, unsigned char >*/,\r
+                0/*addWeighted_gpu<double, unsigned short, signed char >*/,\r
+                0/*addWeighted_gpu<double, unsigned short, unsigned short>*/,\r
+                0/*addWeighted_gpu<double, unsigned short, short >*/,\r
+                0/*addWeighted_gpu<double, unsigned short, int   >*/,\r
+                0/*addWeighted_gpu<double, unsigned short, float >*/,\r
+                0/*addWeighted_gpu<double, unsigned short, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<double, short, unsigned char >*/,\r
+                0/*addWeighted_gpu<double, short, signed char >*/,\r
+                0/*addWeighted_gpu<double, short, unsigned short>*/,\r
+                0/*addWeighted_gpu<double, short, short >*/,\r
+                0/*addWeighted_gpu<double, short, int   >*/,\r
+                0/*addWeighted_gpu<double, short, float >*/,\r
+                0/*addWeighted_gpu<double, short, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<double, int, unsigned char >*/,\r
+                0/*addWeighted_gpu<double, int, signed char >*/,\r
+                0/*addWeighted_gpu<double, int, unsigned short>*/,\r
+                0/*addWeighted_gpu<double, int, short >*/,\r
+                0/*addWeighted_gpu<double, int, int   >*/,\r
+                0/*addWeighted_gpu<double, int, float >*/,\r
+                0/*addWeighted_gpu<double, int, double>*/\r
+            },\r
+            {\r
+                0/*addWeighted_gpu<double, float, unsigned char >*/,\r
+                0/*addWeighted_gpu<double, float, signed char >*/,\r
+                0/*addWeighted_gpu<double, float, unsigned short>*/,\r
+                0/*addWeighted_gpu<double, float, short >*/,\r
+                0/*addWeighted_gpu<double, float, int   >*/,\r
+                0/*addWeighted_gpu<double, float, float >*/,\r
+                0/*addWeighted_gpu<double, float, double>*/\r
+            },\r
+            {\r
+                addWeighted_gpu<double, double, unsigned char >,\r
+                addWeighted_gpu<double, double, signed char >,\r
+                addWeighted_gpu<double, double, unsigned short>,\r
+                addWeighted_gpu<double, double, short >,\r
+                addWeighted_gpu<double, double, int   >,\r
+                addWeighted_gpu<double, double, float >,\r
+                addWeighted_gpu<double, double, double>\r
+            }\r
+        }\r
+    };\r
+\r
+    callers[psrc1->depth()][psrc2->depth()][dst.depth()](psrc1->reshape(1), alpha, psrc2->reshape(1), beta, gamma, dst.reshape(1), StreamAccessor::getStream(stream));\r
+}\r
+\r
 #endif\r

diff --git a/modules/gpu/test/test_arithm.cpp b/modules/gpu/test/test_arithm.cpp
index a8fcccd..31365ca 100644 (file)
--- a/modules/gpu/test/test_arithm.cpp
+++ b/modules/gpu/test/test_arithm.cpp
@@ -1135,7 +1135,7 @@ TEST_P(MinMax, Accuracy)
         return;\r
 \r
     PRINT_PARAM(devInfo);\r
-    PRINT_TYPE(type)\r
+    PRINT_TYPE(type);\r
     PRINT_PARAM(size);\r
 \r
     double minVal, maxVal;\r
@@ -1216,7 +1216,7 @@ TEST_P(MinMaxLoc, Accuracy)
         return;\r
 \r
     PRINT_PARAM(devInfo);\r
-    PRINT_TYPE(type)\r
+    PRINT_TYPE(type);\r
     PRINT_PARAM(size);\r
 \r
     double minVal, maxVal;\r
@@ -1281,7 +1281,7 @@ TEST_P(CountNonZero, Accuracy)
         return;\r
 \r
     PRINT_PARAM(devInfo);\r
-    PRINT_TYPE(type)\r
+    PRINT_TYPE(type);\r
     PRINT_PARAM(size);\r
 \r
     int n;\r
@@ -1333,7 +1333,7 @@ TEST_P(Sum, Accuracy)
         return;\r
 \r
     PRINT_PARAM(devInfo);\r
-    PRINT_TYPE(type)\r
+    PRINT_TYPE(type);\r
     PRINT_PARAM(size);\r
 \r
     cv::Scalar sum;\r
@@ -1385,7 +1385,7 @@ TEST_P(AbsSum, Accuracy)
         return;\r
 \r
     PRINT_PARAM(devInfo);\r
-    PRINT_TYPE(type)\r
+    PRINT_TYPE(type);\r
     PRINT_PARAM(size);\r
 \r
     cv::Scalar sum;\r
@@ -1439,7 +1439,7 @@ TEST_P(SqrSum, Accuracy)
         return;\r
 \r
     PRINT_PARAM(devInfo);\r
-    PRINT_TYPE(type)\r
+    PRINT_TYPE(type);\r
     PRINT_PARAM(size);\r
 \r
     cv::Scalar sum;\r
@@ -1500,7 +1500,7 @@ TEST_P(BitwiseNot, Accuracy)
         return;\r
 \r
     PRINT_PARAM(devInfo);\r
-    PRINT_TYPE(type)\r
+    PRINT_TYPE(type);\r
     PRINT_PARAM(size);\r
 \r
     cv::Mat dst;\r
@@ -1564,7 +1564,7 @@ TEST_P(BitwiseOr, Accuracy)
         return;\r
 \r
     PRINT_PARAM(devInfo);\r
-    PRINT_TYPE(type)\r
+    PRINT_TYPE(type);\r
     PRINT_PARAM(size);\r
 \r
     cv::Mat dst;\r
@@ -1628,7 +1628,7 @@ TEST_P(BitwiseAnd, Accuracy)
         return;\r
 \r
     PRINT_PARAM(devInfo);\r
-    PRINT_TYPE(type)\r
+    PRINT_TYPE(type);\r
     PRINT_PARAM(size);\r
 \r
     cv::Mat dst;\r
@@ -1692,7 +1692,7 @@ TEST_P(BitwiseXor, Accuracy)
         return;\r
 \r
     PRINT_PARAM(devInfo);\r
-    PRINT_TYPE(type)\r
+    PRINT_TYPE(type);\r
     PRINT_PARAM(size);\r
 \r
     cv::Mat dst;\r
@@ -1712,4 +1712,80 @@ INSTANTIATE_TEST_CASE_P(Arithm, BitwiseXor, testing::Combine(
                         testing::ValuesIn(devices()),\r
                         testing::ValuesIn(all_types())));\r
 \r
+//////////////////////////////////////////////////////////////////////////////\r
+// addWeighted\r
+\r
+struct AddWeighted : testing::TestWithParam< std::tr1::tuple<cv::gpu::DeviceInfo, int, int, int> >\r
+{\r
+    cv::gpu::DeviceInfo devInfo;\r
+    int type1;\r
+    int type2;\r
+    int dtype;\r
+\r
+    cv::Size size;\r
+    cv::Mat src1;\r
+    cv::Mat src2;\r
+    double alpha;\r
+    double beta;\r
+    double gamma;\r
+\r
+    cv::Mat dst_gold;\r
+\r
+    virtual void SetUp() \r
+    {\r
+        devInfo = std::tr1::get<0>(GetParam());\r
+        type1 = std::tr1::get<1>(GetParam());\r
+        type2 = std::tr1::get<2>(GetParam());\r
+        dtype = std::tr1::get<3>(GetParam());\r
+\r
+        cv::gpu::setDevice(devInfo.deviceID());\r
+\r
+        cv::RNG& rng = cvtest::TS::ptr()->get_rng();\r
+\r
+        size = cv::Size(rng.uniform(100, 200), rng.uniform(100, 200));\r
+\r
+        src1 = cvtest::randomMat(rng, size, type1, 0.0, 255.0, false);\r
+        src2 = cvtest::randomMat(rng, size, type2, 0.0, 255.0, false);\r
+\r
+        alpha = rng.uniform(-10.0, 10.0);\r
+        beta = rng.uniform(-10.0, 10.0);\r
+        gamma = rng.uniform(-10.0, 10.0);\r
+\r
+        cv::addWeighted(src1, alpha, src2, beta, gamma, dst_gold, dtype);\r
+    }\r
+};\r
+\r
+TEST_P(AddWeighted, Accuracy) \r
+{\r
+    if ((src1.depth() == CV_64F || src2.depth() == CV_64F || dst_gold.depth() == CV_64F) && !supportFeature(devInfo, cv::gpu::NATIVE_DOUBLE))\r
+        return;\r
+\r
+    PRINT_PARAM(devInfo);\r
+    PRINT_TYPE(type1);\r
+    PRINT_TYPE(type2);\r
+    PRINT_TYPE(dtype);\r
+    PRINT_PARAM(size);\r
+    PRINT_PARAM(alpha);\r
+    PRINT_PARAM(beta);\r
+    PRINT_PARAM(gamma);\r
+\r
+    cv::Mat dst;\r
+    \r
+    ASSERT_NO_THROW(\r
+        cv::gpu::GpuMat dev_dst;\r
+\r
+        cv::gpu::addWeighted(cv::gpu::GpuMat(src1), alpha, cv::gpu::GpuMat(src2), beta, gamma, dev_dst, dtype);\r
+\r
+        dev_dst.download(dst);\r
+    );\r
+\r
+    EXPECT_MAT_NEAR(dst_gold, dst, dtype < CV_32F ? 1.0 : 1e-12);\r
+}\r
+\r
+INSTANTIATE_TEST_CASE_P(Arithm, AddWeighted, testing::Combine(\r
+                        testing::ValuesIn(devices()),\r
+                        testing::ValuesIn(types(CV_8U, CV_64F, 1, 1)),\r
+                        testing::ValuesIn(types(CV_8U, CV_64F, 1, 1)),\r
+                        testing::ValuesIn(types(CV_8U, CV_64F, 1, 1))));\r
+\r
 #endif // HAVE_CUDA\r