added missing arithm operations to gpu module

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 51aadeb..e57ccd2 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -539,32 +539,41 @@ namespace cv
         //////////////////////////// Per-element operations ////////////////////////////////////\r
 \r
         //! adds one matrix to another (c = a + b)\r
-        //! supports CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1 types\r
-        CV_EXPORTS void add(const GpuMat& a, const GpuMat& b, GpuMat& c, Stream& stream = Stream::Null());\r
+        CV_EXPORTS void add(const GpuMat& a, const GpuMat& b, GpuMat& c, const GpuMat& mask = GpuMat(), int dtype = -1, Stream& stream = Stream::Null());\r
         //! adds scalar to a matrix (c = a + s)\r
-        //! supports CV_32FC1 and CV_32FC2 type\r
-        CV_EXPORTS void add(const GpuMat& a, const Scalar& sc, GpuMat& c, Stream& stream = Stream::Null());\r
+        CV_EXPORTS void add(const GpuMat& a, const Scalar& sc, GpuMat& c, const GpuMat& mask = GpuMat(), int dtype = -1, Stream& stream = Stream::Null());\r
 \r
         //! subtracts one matrix from another (c = a - b)\r
-        //! supports CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1 types\r
-        CV_EXPORTS void subtract(const GpuMat& a, const GpuMat& b, GpuMat& c, Stream& stream = Stream::Null());\r
+        CV_EXPORTS void subtract(const GpuMat& a, const GpuMat& b, GpuMat& c, const GpuMat& mask = GpuMat(), int dtype = -1, Stream& stream = Stream::Null());\r
         //! subtracts scalar from a matrix (c = a - s)\r
-        //! supports CV_32FC1 and CV_32FC2 type\r
-        CV_EXPORTS void subtract(const GpuMat& a, const Scalar& sc, GpuMat& c, Stream& stream = Stream::Null());\r
-\r
-        //! computes element-wise product of the two arrays (c = a * b)\r
-        //! supports CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1 types\r
-        CV_EXPORTS void multiply(const GpuMat& a, const GpuMat& b, GpuMat& c, Stream& stream = Stream::Null());\r
-        //! multiplies matrix to a scalar (c = a * s)\r
-        //! supports CV_32FC1 type\r
-        CV_EXPORTS void multiply(const GpuMat& a, const Scalar& sc, GpuMat& c, Stream& stream = Stream::Null());\r
-\r
-        //! computes element-wise quotient of the two arrays (c = a / b)\r
-        //! supports CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1 types\r
-        CV_EXPORTS void divide(const GpuMat& a, const GpuMat& b, GpuMat& c, Stream& stream = Stream::Null());\r
-        //! computes element-wise quotient of matrix and scalar (c = a / s)\r
-        //! supports CV_32FC1 type\r
-        CV_EXPORTS void divide(const GpuMat& a, const Scalar& sc, GpuMat& c, Stream& stream = Stream::Null());\r
+        CV_EXPORTS void subtract(const GpuMat& a, const Scalar& sc, GpuMat& c, const GpuMat& mask = GpuMat(), int dtype = -1, Stream& stream = Stream::Null());\r
+\r
+        //! computes element-wise weighted product of the two arrays (c = scale * a * b)\r
+        CV_EXPORTS void multiply(const GpuMat& a, const GpuMat& b, GpuMat& c, double scale = 1, int dtype = -1, Stream& stream = Stream::Null());\r
+        //! weighted multiplies matrix to a scalar (c = scale * a * s)\r
+        CV_EXPORTS void multiply(const GpuMat& a, const Scalar& sc, GpuMat& c, double scale = 1, int dtype = -1, Stream& stream = Stream::Null());\r
+\r
+        //! computes element-wise weighted quotient of the two arrays (c = a / b)\r
+        CV_EXPORTS void divide(const GpuMat& a, const GpuMat& b, GpuMat& c, double scale = 1, int dtype = -1, Stream& stream = Stream::Null());\r
+        //! computes element-wise weighted quotient of matrix and scalar (c = a / s)\r
+        CV_EXPORTS void divide(const GpuMat& a, const Scalar& sc, GpuMat& c, double scale = 1, int dtype = -1, Stream& stream = Stream::Null());\r
+        //! computes element-wise weighted reciprocal of an array (dst = scale/src2)\r
+        CV_EXPORTS void divide(double scale, const GpuMat& src2, GpuMat& dst, int dtype = -1, Stream& stream = Stream::Null());\r
+\r
+        //! computes the weighted sum of two arrays (dst = alpha*src1 + beta*src2 + gamma)\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
+        //! adds scaled array to another one (dst = alpha*src1 + src2)\r
+        static inline void scaleAdd(const GpuMat& src1, double alpha, const GpuMat& src2, GpuMat& dst, Stream& stream = Stream::Null())\r
+        {\r
+            addWeighted(src1, alpha, src2, 1.0, 0.0, dst, -1, stream);\r
+        }\r
+\r
+        //! computes element-wise absolute difference of two arrays (c = abs(a - b))\r
+        CV_EXPORTS void absdiff(const GpuMat& a, const GpuMat& b, GpuMat& c, Stream& stream = Stream::Null());\r
+        //! computes element-wise absolute difference of array and scalar (c = abs(a - s))\r
+        CV_EXPORTS void absdiff(const GpuMat& a, const Scalar& s, GpuMat& c, Stream& stream = Stream::Null());\r
 \r
         //! computes exponent of each matrix element (b = e**a)\r
         //! supports only CV_32FC1 type\r
@@ -580,13 +589,6 @@ namespace cv
         //! supports only CV_32FC1 type\r
         CV_EXPORTS void log(const GpuMat& a, GpuMat& b, Stream& stream = Stream::Null());\r
 \r
-        //! computes element-wise absolute difference of two arrays (c = abs(a - b))\r
-        //! supports CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1 types\r
-        CV_EXPORTS void absdiff(const GpuMat& a, const GpuMat& b, GpuMat& c, Stream& stream = Stream::Null());\r
-        //! computes element-wise absolute difference of array and scalar (c = abs(a - s))\r
-        //! supports only CV_32FC1 type\r
-        CV_EXPORTS void absdiff(const GpuMat& a, const Scalar& s, GpuMat& c, Stream& stream = Stream::Null());\r
-\r
         //! compares elements of two arrays (c = a <cmpop> b)\r
         //! supports CV_8UC4, CV_32FC1 types\r
         CV_EXPORTS void compare(const GpuMat& a, const GpuMat& b, GpuMat& c, int cmpop, Stream& stream = Stream::Null());\r
@@ -615,10 +617,6 @@ 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/src/cuda/element_operations.cu b/modules/gpu/src/cuda/element_operations.cu
index 70f3bab..ce9281c 100644 (file)
--- a/modules/gpu/src/cuda/element_operations.cu
+++ b/modules/gpu/src/cuda/element_operations.cu
@@ -49,32 +49,1197 @@
 \r
 namespace cv { namespace gpu { namespace device\r
 {\r
+    //////////////////////////////////////////////////////////////////////////\r
+    // add\r
+\r
+    template <typename T, typename D> struct Add : binary_function<T, T, D>\r
+    {\r
+        __device__ __forceinline__ D operator ()(T a, T b) const\r
+        {\r
+            return saturate_cast<D>(a + b);\r
+        }\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< Add<ushort, ushort> > : DefaultTransformFunctorTraits< Add<ushort, ushort> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Add<short, short> > : DefaultTransformFunctorTraits< Add<short, short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Add<int, int> > : DefaultTransformFunctorTraits< Add<int, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Add<float, float> > : DefaultTransformFunctorTraits< Add<float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T, typename D> void add_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream)\r
+    {\r
+        if (mask.data)\r
+            transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<D>)dst, mask, Add<T, D>(), stream);\r
+        else\r
+            transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<D>)dst, Add<T, D>(), stream);\r
+    }\r
+\r
+    template void add_gpu<uchar, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<uchar, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<uchar, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<uchar, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<uchar, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<uchar, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<uchar, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<schar, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<ushort, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<ushort, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<ushort, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<ushort, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<ushort, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<ushort, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<ushort, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<short, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<short, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<short, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<short, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<short, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<short, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<short, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<int, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<int, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<int, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<int, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<int, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<int, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<int, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<float, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<float, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<float, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<float, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<float, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<float, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<float, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<double, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<double, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<double, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<double, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<double, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<double, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<double, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    template <typename T, typename D> struct AddScalar : unary_function<T, D>\r
+    {\r
+        AddScalar(double val_) : val(val_) {}\r
+        __device__ __forceinline__ D operator ()(T a) const\r
+        {\r
+            return saturate_cast<D>(a + val);\r
+        }\r
+        const double val;\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< AddScalar<ushort, ushort> > : DefaultTransformFunctorTraits< AddScalar<ushort, ushort>  >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddScalar<short, short> > : DefaultTransformFunctorTraits< AddScalar<short, short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddScalar<int, int> > : DefaultTransformFunctorTraits< AddScalar<int, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AddScalar<float, float> > : DefaultTransformFunctorTraits< AddScalar<float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T, typename D> void add_gpu(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall( cudaSetDoubleForDevice(&val) );\r
+        AddScalar<T, D> op(val);\r
+        if (mask.data)\r
+            transform((DevMem2D_<T>)src1, (DevMem2D_<D>)dst, mask, op, stream);\r
+        else\r
+            transform((DevMem2D_<T>)src1, (DevMem2D_<D>)dst, op, stream);\r
+    }\r
+\r
+    template void add_gpu<uchar, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<uchar, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<uchar, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<uchar, short >(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<uchar, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<uchar, float >(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<uchar, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<schar, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<schar, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<ushort, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<ushort, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<ushort, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<ushort, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<ushort, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<ushort, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<ushort, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<short, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<short, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<short, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<short, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<short, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<short, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<short, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<int, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<int, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<int, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<int, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<int, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<int, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<int, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<float, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<float, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<float, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<float, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<float, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<float, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<float, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void add_gpu<double, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<double, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<double, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<double, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<double, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void add_gpu<double, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void add_gpu<double, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //////////////////////////////////////////////////////////////////////////\r
+    // subtract\r
+\r
+    template <typename T, typename D> struct Subtract : binary_function<T, T, D>\r
+    {\r
+        __device__ __forceinline__ D operator ()(T a, T b) const\r
+        {\r
+            return saturate_cast<D>(a - b);\r
+        }\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< Subtract<ushort, ushort> > : DefaultTransformFunctorTraits< Subtract<ushort, ushort> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Subtract<short, short> > : DefaultTransformFunctorTraits< Subtract<short, short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Subtract<int, int> > : DefaultTransformFunctorTraits< Subtract<int, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Subtract<float, float> > : DefaultTransformFunctorTraits< Subtract<float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T, typename D> void subtract_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream)\r
+    {\r
+        if (mask.data)\r
+            transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<D>)dst, mask, Subtract<T, D>(), stream);\r
+        else\r
+            transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<D>)dst, Subtract<T, D>(), stream);\r
+    }\r
+\r
+    template void subtract_gpu<uchar, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<uchar, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<uchar, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<uchar, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<uchar, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<uchar, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<uchar, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<schar, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<ushort, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<ushort, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<ushort, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<ushort, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<ushort, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<ushort, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<ushort, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<short, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<short, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<short, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<short, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<short, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<short, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<short, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<int, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<int, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<int, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<int, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<int, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<int, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<int, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<float, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<float, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<float, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<float, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<float, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<float, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<float, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<double, uchar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<double, schar>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<double, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<double, short>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<double, int>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<double, float>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<double, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    template <typename T, typename D> struct SubtractScalar : unary_function<T, D>\r
+    {\r
+        SubtractScalar(double val_) : val(val_) {}\r
+        __device__ __forceinline__ D operator ()(T a) const\r
+        {\r
+            return saturate_cast<D>(a - val);\r
+        }\r
+        const double val;\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< SubtractScalar<ushort, ushort> > : DefaultTransformFunctorTraits< SubtractScalar<ushort, ushort>  >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< SubtractScalar<short, short> > : DefaultTransformFunctorTraits< SubtractScalar<short, short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< SubtractScalar<int, int> > : DefaultTransformFunctorTraits< SubtractScalar<int, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< SubtractScalar<float, float> > : DefaultTransformFunctorTraits< SubtractScalar<float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T, typename D> void subtract_gpu(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall( cudaSetDoubleForDevice(&val) );\r
+        SubtractScalar<T, D> op(val);\r
+        if (mask.data)\r
+            transform((DevMem2D_<T>)src1, (DevMem2D_<D>)dst, mask, op, stream);\r
+        else\r
+            transform((DevMem2D_<T>)src1, (DevMem2D_<D>)dst, op, stream);\r
+    }\r
+\r
+    template void subtract_gpu<uchar, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<uchar, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<uchar, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<uchar, short >(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<uchar, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<uchar, float >(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<uchar, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<schar, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<schar, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<ushort, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<ushort, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<ushort, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<ushort, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<ushort, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<ushort, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<ushort, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<short, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<short, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<short, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<short, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<short, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<short, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<short, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<int, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<int, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<int, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<int, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<int, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<int, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<int, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<float, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<float, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<float, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<float, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<float, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<float, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<float, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //template void subtract_gpu<double, uchar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<double, schar>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<double, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<double, short>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<double, int>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    //template void subtract_gpu<double, float>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+    template void subtract_gpu<double, double>(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    //////////////////////////////////////////////////////////////////////////\r
+    // multiply\r
+\r
+    struct multiply_8uc4_32f : binary_function<uint, float, uint>\r
+    {\r
+        __device__ __forceinline__ uint operator ()(uint a, float b) const\r
+        {\r
+            uint res = 0;\r
+\r
+            res |= (saturate_cast<uchar>((0xffu & (a      )) * b)      );\r
+            res |= (saturate_cast<uchar>((0xffu & (a >>  8)) * b) <<  8);\r
+            res |= (saturate_cast<uchar>((0xffu & (a >> 16)) * b) << 16);\r
+            res |= (saturate_cast<uchar>((0xffu & (a >> 24)) * b) << 24);\r
+\r
+            return res;\r
+        }\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits<multiply_8uc4_32f> : DefaultTransformFunctorTraits<multiply_8uc4_32f>\r
+    {\r
+        enum { smart_block_dim_x = 8 };\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 8 };\r
+    };\r
+\r
+    void multiply_gpu(const DevMem2D_<uchar4>& src1, const DevMem2Df& src2, const DevMem2D_<uchar4>& dst, cudaStream_t stream)\r
+    {\r
+        transform(static_cast< DevMem2D_<uint> >(src1), src2, static_cast< DevMem2D_<uint> >(dst), multiply_8uc4_32f(), stream);\r
+    }\r
+\r
+    struct multiply_16sc4_32f : binary_function<short4, float, short4>\r
+    {\r
+        __device__ __forceinline__ short4 operator ()(short4 a, float b) const\r
+        {\r
+            return make_short4(saturate_cast<short>(a.x * b), saturate_cast<short>(a.y * b),\r
+                               saturate_cast<short>(a.z * b), saturate_cast<short>(a.w * b));\r
+        }\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits<multiply_16sc4_32f> : DefaultTransformFunctorTraits<multiply_16sc4_32f>\r
+    {\r
+        enum { smart_block_dim_x = 8 };\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 8 };\r
+    };\r
+\r
+    void multiply_gpu(const DevMem2D_<short4>& src1, const DevMem2Df& src2, const DevMem2D_<short4>& dst, cudaStream_t stream)\r
+    {\r
+        transform(static_cast< DevMem2D_<short4> >(src1), src2, \r
+                  static_cast< DevMem2D_<short4> >(dst), multiply_16sc4_32f(), stream);\r
+    }\r
+\r
+    template <typename T, typename D> struct Multiply : binary_function<T, T, D>\r
+    {\r
+        Multiply(double scale_) : scale(scale_) {}\r
+        __device__ __forceinline__ D operator ()(T a, T b) const\r
+        {\r
+            return saturate_cast<D>(scale * a * b);\r
+        }\r
+        const double scale;\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< Multiply<ushort, ushort> > : DefaultTransformFunctorTraits< Multiply<ushort, ushort> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Multiply<short, short> > : DefaultTransformFunctorTraits< Multiply<short, short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Multiply<int, int> > : DefaultTransformFunctorTraits< Multiply<int, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Multiply<float, float> > : DefaultTransformFunctorTraits< Multiply<float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T, typename D> void multiply_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall( cudaSetDoubleForDevice(&scale) );\r
+        Multiply<T, D> op(scale);\r
+        transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<D>)dst, op, stream);\r
+    }\r
+\r
+    template void multiply_gpu<uchar, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<uchar, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<uchar, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<uchar, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<uchar, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<uchar, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<uchar, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<schar, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<ushort, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<ushort, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<ushort, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<ushort, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<ushort, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<ushort, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<ushort, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<short, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<short, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<short, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<short, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<short, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<short, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<short, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<int, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<int, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<int, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<int, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<int, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<int, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<int, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<float, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<float, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<float, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<float, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<float, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<float, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<float, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<double, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<double, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<double, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<double, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<double, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<double, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<double, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    \r
+    template <typename T, typename D> struct MultiplyScalar : unary_function<T, D>\r
+    {\r
+        MultiplyScalar(double val_, double scale_) : val(val_), scale(scale_) {}\r
+        __device__ __forceinline__ D operator ()(T a) const\r
+        {\r
+            return saturate_cast<D>(scale * a * val);\r
+        }\r
+        const double val;\r
+        const double scale;\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< MultiplyScalar<ushort, ushort> > : DefaultTransformFunctorTraits< MultiplyScalar<ushort, ushort> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< MultiplyScalar<short, short> > : DefaultTransformFunctorTraits< MultiplyScalar<short, short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< MultiplyScalar<int, int> > : DefaultTransformFunctorTraits< MultiplyScalar<int, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< MultiplyScalar<float, float> > : DefaultTransformFunctorTraits< MultiplyScalar<float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T, typename D> void multiply_gpu(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall( cudaSetDoubleForDevice(&val) );\r
+        cudaSafeCall( cudaSetDoubleForDevice(&scale) );\r
+        MultiplyScalar<T, D> op(val, scale);\r
+        transform((DevMem2D_<T>)src1, (DevMem2D_<D>)dst, op, stream);\r
+    }\r
+\r
+    template void multiply_gpu<uchar, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<uchar, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<uchar, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<uchar, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<uchar, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<uchar, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<uchar, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<schar, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<schar, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<ushort, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<ushort, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<ushort, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<ushort, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<ushort, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<ushort, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<ushort, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<short, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<short, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<short, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<short, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<short, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<short, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<short, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<int, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<int, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<int, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<int, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<int, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<int, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<int, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<float, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<float, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<float, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<float, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<float, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<float, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<float, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void multiply_gpu<double, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<double, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<double, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<double, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<double, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void multiply_gpu<double, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void multiply_gpu<double, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //////////////////////////////////////////////////////////////////////////\r
+    // divide\r
+\r
+    struct divide_8uc4_32f : binary_function<uchar4, float, uchar4>\r
+    {\r
+        __device__ __forceinline__ uchar4 operator ()(uchar4 a, float b) const\r
+        {\r
+            return b != 0 ? make_uchar4(saturate_cast<uchar>(a.x / b), saturate_cast<uchar>(a.y / b),\r
+                                        saturate_cast<uchar>(a.z / b), saturate_cast<uchar>(a.w / b)) \r
+                          : make_uchar4(0,0,0,0);\r
+        }\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits<divide_8uc4_32f> : DefaultTransformFunctorTraits<divide_8uc4_32f>\r
+    {\r
+        enum { smart_block_dim_x = 8 };\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 8 };\r
+    };\r
+\r
+    void divide_gpu(const DevMem2D_<uchar4>& src1, const DevMem2Df& src2, const DevMem2D_<uchar4>& dst, cudaStream_t stream)\r
+    {\r
+        transform(static_cast< DevMem2D_<uchar4> >(src1), src2, static_cast< DevMem2D_<uchar4> >(dst), divide_8uc4_32f(), stream);\r
+    }\r
+\r
+\r
+    struct divide_16sc4_32f : binary_function<short4, float, short4>\r
+    {\r
+        __device__ __forceinline__ short4 operator ()(short4 a, float b) const\r
+        {\r
+            return b != 0 ? make_short4(saturate_cast<short>(a.x / b), saturate_cast<uchar>(a.y / b),\r
+                                        saturate_cast<short>(a.z / b), saturate_cast<uchar>(a.w / b))\r
+                          : make_short4(0,0,0,0);\r
+        }\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits<divide_16sc4_32f> : DefaultTransformFunctorTraits<divide_16sc4_32f>\r
+    {\r
+        enum { smart_block_dim_x = 8 };\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 8 };\r
+    };\r
+\r
+    void divide_gpu(const DevMem2D_<short4>& src1, const DevMem2Df& src2, const DevMem2D_<short4>& dst, cudaStream_t stream)\r
+    {\r
+        transform(static_cast< DevMem2D_<short4> >(src1), src2, static_cast< DevMem2D_<short4> >(dst), divide_16sc4_32f(), stream);\r
+    }\r
+\r
+    template <typename T, typename D> struct Divide : binary_function<T, T, D>\r
+    {\r
+        Divide(double scale_) : scale(scale_) {}\r
+        __device__ __forceinline__ D operator ()(T a, T b) const\r
+        {\r
+            return b != 0 ? saturate_cast<D>(scale * a / b) : 0;\r
+        }\r
+        const double scale;\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< Divide<ushort, ushort> > : DefaultTransformFunctorTraits< Divide<ushort, ushort> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Divide<short, short> > : DefaultTransformFunctorTraits< Divide<short, short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Divide<int, int> > : DefaultTransformFunctorTraits< Divide<int, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Divide<float, float> > : DefaultTransformFunctorTraits< Divide<float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T, typename D> void divide_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall( cudaSetDoubleForDevice(&scale) );\r
+        Divide<T, D> op(scale);\r
+        transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<D>)dst, op, stream);\r
+    }\r
+\r
+    template void divide_gpu<uchar, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<uchar, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<uchar, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<uchar, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<uchar, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<uchar, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<uchar, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<schar, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<ushort, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<ushort, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<ushort, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<ushort, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<ushort, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<ushort, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<ushort, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<short, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<short, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<short, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<short, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<short, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<short, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<short, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<int, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<int, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<int, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<int, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<int, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<int, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<int, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<float, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<float, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<float, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<float, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<float, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<float, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<float, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<double, uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<double, schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<double, ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<double, short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<double, int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<double, float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<double, double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    \r
+    template <typename T, typename D> struct DivideScalar : unary_function<T, D>\r
+    {\r
+        DivideScalar(double val_, double scale_) : val(val_), scale(scale_) {}\r
+        __device__ __forceinline__ D operator ()(T a) const\r
+        {\r
+            return saturate_cast<D>(scale * a / val);\r
+        }\r
+        const double val;\r
+        const double scale;\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< DivideScalar<ushort, ushort> > : DefaultTransformFunctorTraits< DivideScalar<ushort, ushort> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< DivideScalar<short, short> > : DefaultTransformFunctorTraits< DivideScalar<short, short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< DivideScalar<int, int> > : DefaultTransformFunctorTraits< DivideScalar<int, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< DivideScalar<float, float> > : DefaultTransformFunctorTraits< DivideScalar<float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T, typename D> void divide_gpu(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall( cudaSetDoubleForDevice(&val) );\r
+        cudaSafeCall( cudaSetDoubleForDevice(&scale) );\r
+        DivideScalar<T, D> op(val, scale);\r
+        transform((DevMem2D_<T>)src1, (DevMem2D_<D>)dst, op, stream);\r
+    }\r
+\r
+    template void divide_gpu<uchar, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<uchar, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<uchar, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<uchar, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<uchar, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<uchar, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<uchar, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<schar, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<schar, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<ushort, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<ushort, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<ushort, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<ushort, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<ushort, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<ushort, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<ushort, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<short, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<short, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<short, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<short, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<short, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<short, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<short, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<int, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<int, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<int, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<int, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<int, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<int, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<int, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<float, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<float, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<float, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<float, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<float, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<float, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<float, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<double, uchar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<double, schar >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<double, ushort>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<double, short >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<double, int   >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    //template void divide_gpu<double, float >(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    template void divide_gpu<double, double>(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    template <typename T, typename D> struct Reciprocal : unary_function<T, D>\r
+    {\r
+        Reciprocal(double scale_) : scale(scale_) {}\r
+        __device__ __forceinline__ D operator ()(T a) const\r
+        {\r
+            return a != 0 ? saturate_cast<D>(scale / a) : 0;\r
+        }\r
+        const double scale;\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< Reciprocal<ushort, ushort> > : DefaultTransformFunctorTraits< Reciprocal<ushort, ushort> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Reciprocal<short, short> > : DefaultTransformFunctorTraits< Reciprocal<short, short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Reciprocal<int, int> > : DefaultTransformFunctorTraits< Reciprocal<int, int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Reciprocal<float, float> > : DefaultTransformFunctorTraits< Reciprocal<float, float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T, typename D> void divide_gpu(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall( cudaSetDoubleForDevice(&scalar) );\r
+        Reciprocal<T, D> op(scalar);\r
+        transform((DevMem2D_<T>)src2, (DevMem2D_<D>)dst, op, stream);\r
+    }\r
+\r
+    template void divide_gpu<uchar, uchar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<uchar, schar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<uchar, ushort>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<uchar, short >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<uchar, int   >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<uchar, float >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<uchar, double>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<schar, uchar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<schar, schar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<schar, ushort>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<schar, short >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<schar, int   >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<schar, float >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<schar, double>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<ushort, uchar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<ushort, schar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<ushort, ushort>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<ushort, short >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<ushort, int   >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<ushort, float >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<ushort, double>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<short, uchar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<short, schar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<short, ushort>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<short, short >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<short, int   >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<short, float >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<short, double>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<int, uchar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<int, schar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<int, ushort>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<int, short >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<int, int   >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<int, float >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<int, double>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<float, uchar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<float, schar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<float, ushort>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<float, short >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<float, int   >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<float, float >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<float, double>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+    //template void divide_gpu<double, uchar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<double, schar >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<double, ushort>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<double, short >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<double, int   >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void divide_gpu<double, float >(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void divide_gpu<double, double>(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+    //////////////////////////////////////////////////////////////////////////\r
+    // absdiff\r
+\r
+    template <typename T> struct Absdiff : binary_function<T, T, T>\r
+    {\r
+        static __device__ __forceinline__ int abs(int a)\r
+        {\r
+            return ::abs(a);\r
+        }\r
+        static __device__ __forceinline__ float abs(float a)\r
+        {\r
+            return ::fabsf(a);\r
+        }\r
+        static __device__ __forceinline__ double abs(double a)\r
+        {\r
+            return ::fabs(a);\r
+        }\r
+\r
+        __device__ __forceinline__ T operator ()(T a, T b) const\r
+        {\r
+            return saturate_cast<T>(abs(a - b));\r
+        }\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< Absdiff<ushort> > : DefaultTransformFunctorTraits< Absdiff<ushort> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Absdiff<short> > : DefaultTransformFunctorTraits< Absdiff<short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Absdiff<int> > : DefaultTransformFunctorTraits< Absdiff<int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Absdiff<float> > : DefaultTransformFunctorTraits< Absdiff<float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T> void absdiff_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream)\r
+    {\r
+        transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<T>)dst, Absdiff<T>(), stream);\r
+    }\r
+\r
+    //template void absdiff_gpu<uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void absdiff_gpu<schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void absdiff_gpu<ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void absdiff_gpu<short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void absdiff_gpu<int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    //template void absdiff_gpu<float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void absdiff_gpu<double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+    template <typename T> struct AbsdiffScalar : unary_function<T, T>\r
+    {\r
+        AbsdiffScalar(double val_) : val(val_) {}\r
+        __device__ __forceinline__ T operator ()(T a) const\r
+        {\r
+            return saturate_cast<T>(::fabs(a - val));\r
+        }\r
+        double val;\r
+    };\r
+\r
+    template <> struct TransformFunctorTraits< AbsdiffScalar<ushort> > : DefaultTransformFunctorTraits< AbsdiffScalar<ushort> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AbsdiffScalar<short> > : DefaultTransformFunctorTraits< AbsdiffScalar<short> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AbsdiffScalar<int> > : DefaultTransformFunctorTraits< AbsdiffScalar<int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< AbsdiffScalar<float> > : DefaultTransformFunctorTraits< AbsdiffScalar<float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <typename T> void absdiff_gpu(const DevMem2D& src1, double val, const DevMem2D& dst, cudaStream_t stream)\r
+    {\r
+        cudaSafeCall( cudaSetDoubleForDevice(&val) );\r
+        AbsdiffScalar<T> op(val);\r
+        transform((DevMem2D_<T>)src1, (DevMem2D_<T>)dst, op, stream);\r
+    }\r
+\r
+    template void absdiff_gpu<uchar >(const DevMem2D& src1, double src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void absdiff_gpu<schar >(const DevMem2D& src1, double src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void absdiff_gpu<ushort>(const DevMem2D& src1, double src2, const DevMem2D& dst, cudaStream_t stream);                  \r
+    template void absdiff_gpu<short >(const DevMem2D& src1, double src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void absdiff_gpu<int   >(const DevMem2D& src1, double src2, const DevMem2D& dst, cudaStream_t stream);                  \r
+    //template void absdiff_gpu<float >(const DevMem2D& src1, double src2, const DevMem2D& dst, cudaStream_t stream);                  \r
+    template void absdiff_gpu<double>(const DevMem2D& src1, double src2, const DevMem2D& dst, cudaStream_t stream);\r
+\r
     //////////////////////////////////////////////////////////////////////////////////////\r
     // Compare\r
 \r
+    template <typename T> struct Equal : binary_function<T, T, uchar>\r
+    {\r
+        __device__ __forceinline__ uchar operator()(T src1, T src2) const\r
+        {\r
+            return static_cast<uchar>((src1 == src2) * 255);\r
+        }\r
+    };\r
     template <typename T> struct NotEqual : binary_function<T, T, uchar>\r
     {\r
         __device__ __forceinline__ uchar operator()(T src1, T src2) const\r
         {\r
-            return static_cast<uchar>(static_cast<int>(src1 != src2) * 255);\r
+            return static_cast<uchar>((src1 != src2) * 255);\r
+        }\r
+    };\r
+    template <typename T> struct Less : binary_function<T, T, uchar>\r
+    {\r
+        __device__ __forceinline__ uchar operator()(T src1, T src2) const\r
+        {\r
+            return static_cast<uchar>((src1 < src2) * 255);\r
+        }\r
+    };\r
+    template <typename T> struct LessEqual : binary_function<T, T, uchar>\r
+    {\r
+        __device__ __forceinline__ uchar operator()(T src1, T src2) const\r
+        {\r
+            return static_cast<uchar>((src1 <= src2) * 255);\r
         }\r
     };\r
 \r
-    template <typename T>\r
-    inline void compare_ne(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream)\r
+    template <> struct TransformFunctorTraits< Equal<int> > : DefaultTransformFunctorTraits< Equal<int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Equal<float> > : DefaultTransformFunctorTraits< Equal<float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< NotEqual<int> > : DefaultTransformFunctorTraits< NotEqual<int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< NotEqual<float> > : DefaultTransformFunctorTraits< NotEqual<float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Less<int> > : DefaultTransformFunctorTraits< Less<int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< Less<float> > : DefaultTransformFunctorTraits< Less<float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< LessEqual<int> > : DefaultTransformFunctorTraits< LessEqual<int> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+    template <> struct TransformFunctorTraits< LessEqual<float> > : DefaultTransformFunctorTraits< LessEqual<float> >\r
+    {\r
+        enum { smart_block_dim_y = 8 };\r
+        enum { smart_shift = 4 };\r
+    };\r
+\r
+    template <template <typename> class Op, typename T> void compare(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream)\r
     {\r
-        NotEqual<T> op;\r
+        Op<T> op;\r
         transform(static_cast< DevMem2D_<T> >(src1), static_cast< DevMem2D_<T> >(src2), dst, op, stream);\r
     }\r
 \r
-    void compare_ne_8uc4(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream)\r
+    template <typename T> void compare_eq(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream)\r
+    {\r
+        compare<Equal, T>(src1, src2, dst, stream);\r
+    }\r
+    template <typename T> void compare_ne(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream)\r
     {\r
-        compare_ne<uint>(src1, src2, dst, stream);\r
+        compare<NotEqual, T>(src1, src2, dst, stream);\r
     }\r
-    void compare_ne_32f(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream)\r
+    template <typename T> void compare_lt(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream)\r
     {\r
-        compare_ne<float>(src1, src2, dst, stream);\r
+        compare<Less, T>(src1, src2, dst, stream);\r
     }\r
+    template <typename T> void compare_le(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream)\r
+    {\r
+        compare<LessEqual, T>(src1, src2, dst, stream);\r
+    }\r
+    \r
+    template void compare_eq<uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_eq<schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_eq<ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_eq<short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_eq<int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_eq<float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_eq<double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    \r
+    template void compare_ne<uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_ne<schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_ne<ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_ne<short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_ne<int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_ne<float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_ne<double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    \r
+    template void compare_lt<uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_lt<schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_lt<ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_lt<short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_lt<int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_lt<float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_lt<double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    \r
+    template void compare_le<uchar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_le<schar >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_le<ushort>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_le<short >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_le<int   >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_le<float >(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template void compare_le<double>(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
 \r
 \r
     //////////////////////////////////////////////////////////////////////////\r
@@ -508,21 +1673,6 @@ namespace cv { namespace gpu { namespace device
     template void threshold_gpu<double>(const DevMem2D& src, const DevMem2D& dst, double thresh, double maxVal, int type, cudaStream_t stream);\r
 \r
 \r
-    //////////////////////////////////////////////////////////////////////////\r
-    // subtract\r
-\r
-    template <> struct TransformFunctorTraits< minus<short> > : DefaultTransformFunctorTraits< minus<short> >\r
-    {\r
-        enum { smart_block_dim_y = 8 };\r
-        enum { smart_shift = 4 };\r
-    };\r
-\r
-    template <typename T> void subtractCaller(const DevMem2D src1, const DevMem2D src2, DevMem2D dst, cudaStream_t stream)\r
-    {\r
-        transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<T>)dst, minus<T>(), stream);\r
-    }\r
-\r
-    template void subtractCaller<short>(const DevMem2D src1, const DevMem2D src2, DevMem2D dst, cudaStream_t stream);\r
 \r
 \r
     //////////////////////////////////////////////////////////////////////////\r
@@ -604,151 +1754,7 @@ namespace cv { namespace gpu { namespace device
     template void pow_caller<float>(const DevMem2D& src, float power, DevMem2D dst, cudaStream_t stream);\r
 \r
 \r
-    //////////////////////////////////////////////////////////////////////////\r
-    // divide\r
-\r
-    struct divide_8uc4_32f : binary_function<uchar4, float, uchar4>\r
-    {\r
-        __device__ __forceinline__ uchar4 operator ()(uchar4 a, float b) const\r
-        {\r
-            return make_uchar4(saturate_cast<uchar>(a.x / b), saturate_cast<uchar>(a.y / b),\r
-                               saturate_cast<uchar>(a.z / b), saturate_cast<uchar>(a.w / b));\r
-        }\r
-    };\r
-\r
-    template <> struct TransformFunctorTraits<divide_8uc4_32f> : DefaultTransformFunctorTraits<divide_8uc4_32f>\r
-    {\r
-        enum { smart_block_dim_x = 8 };\r
-        enum { smart_block_dim_y = 8 };\r
-        enum { smart_shift = 8 };\r
-    };\r
-\r
-    void divide_gpu(const DevMem2D_<uchar4>& src1, const DevMem2Df& src2, const DevMem2D_<uchar4>& dst, cudaStream_t stream)\r
-    {\r
-        transform(static_cast< DevMem2D_<uchar4> >(src1), src2, static_cast< DevMem2D_<uchar4> >(dst), divide_8uc4_32f(), stream);\r
-    }\r
-\r
-\r
-    struct divide_16sc4_32f : binary_function<short4, float, short4>\r
-    {\r
-        __device__ __forceinline__ short4 operator ()(short4 a, float b) const\r
-        {\r
-            return make_short4(saturate_cast<short>(a.x / b), saturate_cast<uchar>(a.y / b),\r
-                               saturate_cast<short>(a.z / b), saturate_cast<uchar>(a.w / b));\r
-        }\r
-    };\r
-\r
-    template <> struct TransformFunctorTraits<divide_16sc4_32f> : DefaultTransformFunctorTraits<divide_16sc4_32f>\r
-    {\r
-        enum { smart_block_dim_x = 8 };\r
-        enum { smart_block_dim_y = 8 };\r
-        enum { smart_shift = 8 };\r
-    };\r
-\r
-    void divide_gpu(const DevMem2D_<short4>& src1, const DevMem2Df& src2, const DevMem2D_<short4>& dst, cudaStream_t stream)\r
-    {\r
-        transform(static_cast< DevMem2D_<short4> >(src1), src2, static_cast< DevMem2D_<short4> >(dst), divide_16sc4_32f(), stream);\r
-    }\r
-\r
-\r
-    //////////////////////////////////////////////////////////////////////////\r
-    // multiply\r
-\r
-    template <> struct TransformFunctorTraits< plus<short> > : DefaultTransformFunctorTraits< plus<short> >\r
-    {\r
-        enum { smart_block_dim_y = 8 };\r
-        enum { smart_shift = 4 };\r
-    };\r
-\r
-    template <typename T> void add_gpu(const DevMem2D src1, const DevMem2D src2, DevMem2D dst, cudaStream_t stream)\r
-    {\r
-        transform((DevMem2D_<T>)src1, (DevMem2D_<T>)src2, (DevMem2D_<T>)dst, plus<T>(), stream);\r
-    }\r
-\r
-    template void add_gpu<short>(const DevMem2D src1, const DevMem2D src2, DevMem2D dst, cudaStream_t stream);\r
-\r
-    //////////////////////////////////////////////////////////////////////////\r
-    // multiply\r
-\r
-    struct multiply_8uc4_32f : binary_function<uint, float, uint>\r
-    {\r
-        __device__ __forceinline__ uint operator ()(uint a, float b) const\r
-        {\r
-            uint res = 0;\r
-\r
-            res |= (saturate_cast<uchar>((0xffu & (a      )) * b)      );\r
-            res |= (saturate_cast<uchar>((0xffu & (a >>  8)) * b) <<  8);\r
-            res |= (saturate_cast<uchar>((0xffu & (a >> 16)) * b) << 16);\r
-            res |= (saturate_cast<uchar>((0xffu & (a >> 24)) * b) << 24);\r
-\r
-            return res;\r
-        }\r
-    };\r
-\r
-    template <> struct TransformFunctorTraits<multiply_8uc4_32f> : DefaultTransformFunctorTraits<multiply_8uc4_32f>\r
-    {\r
-        enum { smart_block_dim_x = 8 };\r
-        enum { smart_block_dim_y = 8 };\r
-        enum { smart_shift = 8 };\r
-    };\r
-\r
-    void multiply_gpu(const DevMem2D_<uchar4>& src1, const DevMem2Df& src2, const DevMem2D_<uchar4>& dst, cudaStream_t stream)\r
-    {\r
-        transform(static_cast< DevMem2D_<uint> >(src1), src2, static_cast< DevMem2D_<uint> >(dst), multiply_8uc4_32f(), stream);\r
-    }\r
-\r
-    struct multiply_16sc4_32f : binary_function<short4, float, short4>\r
-    {\r
-        __device__ __forceinline__ short4 operator ()(short4 a, float b) const\r
-        {\r
-            return make_short4(saturate_cast<short>(a.x * b), saturate_cast<short>(a.y * b),\r
-                               saturate_cast<short>(a.z * b), saturate_cast<short>(a.w * b));\r
-        }\r
-    };\r
-\r
-    template <> struct TransformFunctorTraits<multiply_16sc4_32f> : DefaultTransformFunctorTraits<multiply_16sc4_32f>\r
-    {\r
-        enum { smart_block_dim_x = 8 };\r
-        enum { smart_block_dim_y = 8 };\r
-        enum { smart_shift = 8 };\r
-    };\r
-\r
-    void multiply_gpu(const DevMem2D_<short4>& src1, const DevMem2Df& src2, const DevMem2D_<short4>& dst, cudaStream_t stream)\r
-    {\r
-        transform(static_cast< DevMem2D_<short4> >(src1), src2, \r
-                  static_cast< DevMem2D_<short4> >(dst), multiply_16sc4_32f(), stream);\r
-    }\r
-\r
-\r
-    //////////////////////////////////////////////////////////////////////////\r
-    // multiply (by scalar)\r
-\r
-    template <typename T, typename D, typename S> struct MultiplyScalar : unary_function<T, D>\r
-    {\r
-        __host__ __device__ __forceinline__ MultiplyScalar(typename TypeTraits<S>::ParameterType scale_) : scale(scale_) {}\r
-\r
-        __device__ __forceinline__ D operator ()(typename TypeTraits<T>::ParameterType a) const\r
-        {\r
-            return saturate_cast<D>(a * scale);\r
-        }\r
-\r
-        const S scale;\r
-    };\r
-\r
-    template <> struct TransformFunctorTraits< MultiplyScalar<uchar, uchar, float> > : DefaultTransformFunctorTraits< MultiplyScalar<uchar, uchar, float> >\r
-    {\r
-        enum { smart_block_dim_y = 8 };\r
-        enum { smart_shift = 8 };\r
-    };\r
-\r
-    template <typename T, typename D>\r
-    void multiplyScalar_gpu(const DevMem2D& src, float scale, const DevMem2D& dst, cudaStream_t stream)\r
-    {\r
-        transform(static_cast< DevMem2D_<T> >(src), static_cast< DevMem2D_<D> >(dst), MultiplyScalar<T, D, float>(scale), stream);\r
-    }\r
-\r
-    template void multiplyScalar_gpu<uchar, uchar>(const DevMem2D& src, float scale, const DevMem2D& dst, cudaStream_t stream);\r
-\r
+    \r
 \r
     //////////////////////////////////////////////////////////////////////////\r
     // addWeighted\r

diff --git a/modules/gpu/src/element_operations.cpp b/modules/gpu/src/element_operations.cpp
index 1173803..b816de2 100644 (file)
--- a/modules/gpu/src/element_operations.cpp
+++ b/modules/gpu/src/element_operations.cpp
@@ -47,14 +47,15 @@ using namespace cv::gpu;
 \r
 #if !defined (HAVE_CUDA)\r
 \r
-void cv::gpu::add(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
-void cv::gpu::add(const GpuMat&, const Scalar&, GpuMat&, Stream&) { throw_nogpu(); }\r
-void cv::gpu::subtract(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
-void cv::gpu::subtract(const GpuMat&, const Scalar&, GpuMat&, Stream&) { throw_nogpu(); }\r
-void cv::gpu::multiply(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
-void cv::gpu::multiply(const GpuMat&, const Scalar&, GpuMat&, Stream&) { throw_nogpu(); }\r
-void cv::gpu::divide(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
-void cv::gpu::divide(const GpuMat&, const Scalar&, GpuMat&, Stream&) { throw_nogpu(); }\r
+void cv::gpu::add(const GpuMat&, const GpuMat&, GpuMat&, const GpuMat&, int, Stream&) { throw_nogpu(); }\r
+void cv::gpu::add(const GpuMat&, const Scalar&, GpuMat&, const GpuMat&, int, Stream&) { throw_nogpu(); }\r
+void cv::gpu::subtract(const GpuMat&, const GpuMat&, GpuMat&, const GpuMat&, int, Stream&) { throw_nogpu(); }\r
+void cv::gpu::subtract(const GpuMat&, const Scalar&, GpuMat&, const GpuMat&, int, Stream&) { throw_nogpu(); }\r
+void cv::gpu::multiply(const GpuMat&, const GpuMat&, GpuMat&, double, int, Stream&) { throw_nogpu(); }\r
+void cv::gpu::multiply(const GpuMat&, const Scalar&, GpuMat&, double, int, Stream&) { throw_nogpu(); }\r
+void cv::gpu::divide(const GpuMat&, const GpuMat&, GpuMat&, double, int, Stream&) { throw_nogpu(); }\r
+void cv::gpu::divide(const GpuMat&, const Scalar&, GpuMat&, double, int, Stream&) { throw_nogpu(); }\r
+void cv::gpu::divide(double, const GpuMat&, GpuMat&, int, Stream&) { throw_nogpu(); }\r
 void cv::gpu::absdiff(const GpuMat&, const GpuMat&, GpuMat&, Stream&) { throw_nogpu(); }\r
 void cv::gpu::absdiff(const GpuMat&, const Scalar&, GpuMat&, Stream&) { throw_nogpu(); }\r
 void cv::gpu::compare(const GpuMat&, const GpuMat&, GpuMat&, int, Stream&) { throw_nogpu(); }\r
@@ -85,151 +86,305 @@ namespace
                          npp_arithm_8u_t npp_func_8uc1, npp_arithm_8u_t npp_func_8uc4,\r
                          npp_arithm_32s_t npp_func_32sc1, npp_arithm_32f_t npp_func_32fc1, cudaStream_t stream)\r
     {\r
-        CV_DbgAssert(src1.size() == src2.size() && src1.type() == src2.type());\r
-        CV_Assert(src1.type() == CV_8UC1 || src1.type() == CV_8UC4 || src1.type() == CV_32SC1 || src1.type() == CV_32FC1);\r
-        dst.create( src1.size(), src1.type() );\r
-\r
         NppiSize sz;\r
-        sz.width  = src1.cols;\r
+        sz.width  = src1.cols * src1.channels();\r
         sz.height = src1.rows;\r
 \r
         NppStreamHandler h(stream);\r
 \r
-        switch (src1.type())\r
+        if (src1.depth() == CV_8U && (src1.cols * src1.channels()) % 4 == 0)\r
         {\r
-        case CV_8UC1:\r
-            nppSafeCall( npp_func_8uc1(src1.ptr<Npp8u>(), static_cast<int>(src1.step), src2.ptr<Npp8u>(), static_cast<int>(src2.step), \r
-                dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz, 0) );\r
-            break;\r
-        case CV_8UC4:\r
+            sz.width /= 4;\r
+\r
             nppSafeCall( npp_func_8uc4(src1.ptr<Npp8u>(), static_cast<int>(src1.step), src2.ptr<Npp8u>(), static_cast<int>(src2.step), \r
                 dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz, 0) );\r
-            break;\r
-        case CV_32SC1:\r
+        }\r
+        else if (src1.depth() == CV_8U)\r
+        {\r
+            nppSafeCall( npp_func_8uc1(src1.ptr<Npp8u>(), static_cast<int>(src1.step), src2.ptr<Npp8u>(), static_cast<int>(src2.step), \r
+                dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz, 0) );\r
+        }\r
+        else if (src1.depth() == CV_32S)\r
+        {\r
             nppSafeCall( npp_func_32sc1(src1.ptr<Npp32s>(), static_cast<int>(src1.step), src2.ptr<Npp32s>(), static_cast<int>(src2.step), \r
                 dst.ptr<Npp32s>(), static_cast<int>(dst.step), sz) );\r
-            break;\r
-        case CV_32FC1:\r
+        }\r
+        else if (src1.depth() == CV_32F)\r
+        {\r
             nppSafeCall( npp_func_32fc1(src1.ptr<Npp32f>(), static_cast<int>(src1.step), src2.ptr<Npp32f>(), static_cast<int>(src2.step), \r
                 dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz) );\r
-            break;\r
-        default:\r
-            CV_Assert(!"Unsupported source type");\r
         }\r
 \r
         if (stream == 0)\r
             cudaSafeCall( cudaDeviceSynchronize() );\r
     }\r
+}\r
 \r
-    template<int SCN> struct NppArithmScalarFunc;\r
-    template<> struct NppArithmScalarFunc<1>\r
-    {\r
-        typedef NppStatus (*func_ptr)(const Npp32f *pSrc, int nSrcStep, Npp32f nValue, Npp32f *pDst, int nDstStep, NppiSize oSizeROI);\r
-    };\r
-    template<> struct NppArithmScalarFunc<2>\r
+////////////////////////////////////////////////////////////////////////\r
+// add\r
+\r
+namespace cv { namespace gpu { namespace device\r
+{\r
+    template <typename T, typename D> \r
+    void add_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    template <typename T, typename D> \r
+    void add_gpu(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+}}}\r
+\r
+void cv::gpu::add(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, int dtype, Stream& s)\r
+{\r
+    using namespace cv::gpu::device;\r
+\r
+    typedef void (*func_t)(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    static const func_t funcs[7][7] = \r
     {\r
-        typedef NppStatus (*func_ptr)(const Npp32fc *pSrc, int nSrcStep, Npp32fc nValue, Npp32fc *pDst, int nDstStep, NppiSize oSizeROI);\r
+        {add_gpu<unsigned char, unsigned char>, 0/*add_gpu<unsigned char, signed char>*/, add_gpu<unsigned char, unsigned short>, add_gpu<unsigned char, short>, add_gpu<unsigned char, int>, add_gpu<unsigned char, float>, add_gpu<unsigned char, double>},\r
+        {0/*add_gpu<signed char, unsigned char>*/, 0/*add_gpu<signed char, signed char>*/, 0/*add_gpu<signed char, unsigned short>*/, 0/*add_gpu<signed char, short>*/, 0/*add_gpu<signed char, int>*/, 0/*add_gpu<signed char, float>*/, 0/*add_gpu<signed char, double>*/},\r
+        {0/*add_gpu<unsigned short, unsigned char>*/, 0/*add_gpu<unsigned short, signed char>*/, add_gpu<unsigned short, unsigned short>, 0/*add_gpu<unsigned short, short>*/, add_gpu<unsigned short, int>, add_gpu<unsigned short, float>, add_gpu<unsigned short, double>},\r
+        {0/*add_gpu<short, unsigned char>*/, 0/*add_gpu<short, signed char>*/, 0/*add_gpu<short, unsigned short>*/, add_gpu<short, short>, add_gpu<short, int>, add_gpu<short, float>, add_gpu<short, double>},\r
+        {0/*add_gpu<int, unsigned char>*/, 0/*add_gpu<int, signed char>*/, 0/*add_gpu<int, unsigned short>*/, 0/*add_gpu<int, short>*/, add_gpu<int, int>, add_gpu<int, float>, add_gpu<int, double>},\r
+        {0/*add_gpu<float, unsigned char>*/, 0/*add_gpu<float, signed char>*/, 0/*add_gpu<float, unsigned short>*/, 0/*add_gpu<float, short>*/, 0/*add_gpu<float, int>*/, add_gpu<float, float>, add_gpu<float, double>},\r
+        {0/*add_gpu<double, unsigned char>*/, 0/*add_gpu<double, signed char>*/, 0/*add_gpu<double, unsigned short>*/, 0/*add_gpu<double, short>*/, 0/*add_gpu<double, int>*/, 0/*add_gpu<double, float>*/, add_gpu<double, double>}\r
     };\r
 \r
-    template<int SCN, typename NppArithmScalarFunc<SCN>::func_ptr func> struct NppArithmScalar;\r
+    CV_Assert(src1.type() == src2.type() && src1.size() == src2.size());\r
+    CV_Assert(mask.empty() || (src1.channels() == 1 && mask.size() == src1.size() && mask.type() == CV_8U));\r
+\r
+    if (dtype < 0)\r
+        dtype = src1.depth();\r
+\r
+    dst.create(src1.size(), CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src1.channels()));\r
+\r
+    cudaStream_t stream = StreamAccessor::getStream(s);\r
 \r
-    template<typename NppArithmScalarFunc<1>::func_ptr func> struct NppArithmScalar<1, func>\r
+    if (mask.empty() && dst.type() == src1.type() && (src1.depth() == CV_8U || src1.depth() == CV_32S || src1.depth() == CV_32F))\r
     {\r
-        static void calc(const GpuMat& src, const Scalar& sc, GpuMat& dst, cudaStream_t stream)\r
-        {\r
-            dst.create(src.size(), src.type());\r
+        nppArithmCaller(src1, src2, dst, nppiAdd_8u_C1RSfs, nppiAdd_8u_C4RSfs, nppiAdd_32s_C1R, nppiAdd_32f_C1R, stream);\r
+        return;\r
+    }\r
 \r
-            NppiSize sz;\r
-            sz.width  = src.cols;\r
-            sz.height = src.rows;\r
+    const func_t func = funcs[src1.depth()][dst.depth()];\r
+    CV_Assert(func != 0);\r
+\r
+    func(src1.reshape(1), src2.reshape(1), dst.reshape(1), mask, stream);\r
+}\r
 \r
-            NppStreamHandler h(stream);\r
+void cv::gpu::add(const GpuMat& src, const Scalar& sc, GpuMat& dst, const GpuMat& mask, int dtype, Stream& s)\r
+{\r
+    using namespace cv::gpu::device;\r
 \r
-            nppSafeCall( func(src.ptr<Npp32f>(), static_cast<int>(src.step), (Npp32f)sc[0], dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz) );\r
+    typedef void (*func_t)(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
 \r
-            if (stream == 0)\r
-                cudaSafeCall( cudaDeviceSynchronize() );\r
-        }\r
+    static const func_t funcs[7][7] = \r
+    {\r
+        {add_gpu<unsigned char, unsigned char>, 0/*add_gpu<unsigned char, signed char>*/, add_gpu<unsigned char, unsigned short>, add_gpu<unsigned char, short>, add_gpu<unsigned char, int>, add_gpu<unsigned char, float>, add_gpu<unsigned char, double>},\r
+        {0/*add_gpu<signed char, unsigned char>*/, 0/*add_gpu<signed char, signed char>*/, 0/*add_gpu<signed char, unsigned short>*/, 0/*add_gpu<signed char, short>*/, 0/*add_gpu<signed char, int>*/, 0/*add_gpu<signed char, float>*/, 0/*add_gpu<signed char, double>*/},\r
+        {0/*add_gpu<unsigned short, unsigned char>*/, 0/*add_gpu<unsigned short, signed char>*/, add_gpu<unsigned short, unsigned short>, 0/*add_gpu<unsigned short, short>*/, add_gpu<unsigned short, int>, add_gpu<unsigned short, float>, add_gpu<unsigned short, double>},\r
+        {0/*add_gpu<short, unsigned char>*/, 0/*add_gpu<short, signed char>*/, 0/*add_gpu<short, unsigned short>*/, add_gpu<short, short>, add_gpu<short, int>, add_gpu<short, float>, add_gpu<short, double>},\r
+        {0/*add_gpu<int, unsigned char>*/, 0/*add_gpu<int, signed char>*/, 0/*add_gpu<int, unsigned short>*/, 0/*add_gpu<int, short>*/, add_gpu<int, int>, add_gpu<int, float>, add_gpu<int, double>},\r
+        {0/*add_gpu<float, unsigned char>*/, 0/*add_gpu<float, signed char>*/, 0/*add_gpu<float, unsigned short>*/, 0/*add_gpu<float, short>*/, 0/*add_gpu<float, int>*/, add_gpu<float, float>, add_gpu<float, double>},\r
+        {0/*add_gpu<double, unsigned char>*/, 0/*add_gpu<double, signed char>*/, 0/*add_gpu<double, unsigned short>*/, 0/*add_gpu<double, short>*/, 0/*add_gpu<double, int>*/, 0/*add_gpu<double, float>*/, add_gpu<double, double>}\r
     };\r
-    template<typename NppArithmScalarFunc<2>::func_ptr func> struct NppArithmScalar<2, func>\r
+\r
+    CV_Assert(src.channels() == 1 || src.type() == CV_32FC2);\r
+    CV_Assert(mask.empty() || (src.channels() == 1 && mask.size() == src.size() && mask.type() == CV_8U));\r
+\r
+    if (dtype < 0)\r
+        dtype = src.depth();\r
+\r
+    dst.create(src.size(), CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src.channels()));\r
+\r
+    cudaStream_t stream = StreamAccessor::getStream(s);\r
+\r
+    if (mask.empty() && dst.type() == src.type() && src.depth() == CV_32F)\r
     {\r
-        static void calc(const GpuMat& src, const Scalar& sc, GpuMat& dst, cudaStream_t stream)\r
-        {\r
-            dst.create(src.size(), src.type());\r
+        NppiSize sz;\r
+        sz.width  = src.cols;\r
+        sz.height = src.rows;\r
 \r
-            NppiSize sz;\r
-            sz.width  = src.cols;\r
-            sz.height = src.rows;\r
+        NppStreamHandler h(stream);\r
 \r
+        if (src.type() == CV_32FC1)\r
+        {\r
+            nppSafeCall( nppiAddC_32f_C1R(src.ptr<Npp32f>(), static_cast<int>(src.step), static_cast<Npp32f>(sc.val[0]), \r
+                dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz) );\r
+        }\r
+        else\r
+        {\r
             Npp32fc nValue;\r
-            nValue.re = (Npp32f)sc[0];\r
-            nValue.im = (Npp32f)sc[1];\r
+            nValue.re = static_cast<Npp32f>(sc.val[0]);\r
+            nValue.im = static_cast<Npp32f>(sc.val[1]);\r
+            nppSafeCall( nppiAddC_32fc_C1R(src.ptr<Npp32fc>(), static_cast<int>(src.step), nValue, \r
+                dst.ptr<Npp32fc>(), static_cast<int>(dst.step), sz) );\r
+        }\r
 \r
-            NppStreamHandler h(stream);\r
+        if (stream == 0)\r
+            cudaSafeCall( cudaDeviceSynchronize() );\r
 \r
-            nppSafeCall( func(src.ptr<Npp32fc>(), static_cast<int>(src.step), nValue, dst.ptr<Npp32fc>(), static_cast<int>(dst.step), sz) );\r
+        return;\r
+    }\r
 \r
-            if (stream == 0)\r
-                cudaSafeCall( cudaDeviceSynchronize() );\r
-        }\r
-    };\r
+    const func_t func = funcs[src.depth()][dst.depth()];\r
+    CV_Assert(func != 0);\r
+\r
+    func(src, sc.val[0], dst, mask, stream);\r
 }\r
 \r
+////////////////////////////////////////////////////////////////////////\r
+// subtract\r
+\r
 namespace cv { namespace gpu { namespace device\r
 {\r
-    template <typename T>\r
-    void add_gpu(const DevMem2D src1, const DevMem2D src2, DevMem2D dst, cudaStream_t stream);\r
+    template <typename T, typename D> \r
+    void subtract_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    template <typename T, typename D> \r
+    void subtract_gpu(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
 }}}\r
 \r
-void cv::gpu::add(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)\r
+void cv::gpu::subtract(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, const GpuMat& mask, int dtype, Stream& s)\r
 {\r
-    if (src1.depth() == CV_16S && src2.depth() == CV_16S)\r
+    using namespace cv::gpu::device;\r
+\r
+    typedef void (*func_t)(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    static const func_t funcs[7][7] = \r
     {\r
-        CV_Assert(src1.size() == src2.size());\r
-        dst.create(src1.size(), src1.type());\r
-        device::add_gpu<short>(src1.reshape(1), src2.reshape(1), dst.reshape(1), StreamAccessor::getStream(stream));\r
+        {subtract_gpu<unsigned char, unsigned char>, 0/*subtract_gpu<unsigned char, signed char>*/, subtract_gpu<unsigned char, unsigned short>, subtract_gpu<unsigned char, short>, subtract_gpu<unsigned char, int>, subtract_gpu<unsigned char, float>, subtract_gpu<unsigned char, double>},\r
+        {0/*subtract_gpu<signed char, unsigned char>*/, 0/*subtract_gpu<signed char, signed char>*/, 0/*subtract_gpu<signed char, unsigned short>*/, 0/*subtract_gpu<signed char, short>*/, 0/*subtract_gpu<signed char, int>*/, 0/*subtract_gpu<signed char, float>*/, 0/*subtract_gpu<signed char, double>*/},\r
+        {0/*subtract_gpu<unsigned short, unsigned char>*/, 0/*subtract_gpu<unsigned short, signed char>*/, subtract_gpu<unsigned short, unsigned short>, 0/*subtract_gpu<unsigned short, short>*/, subtract_gpu<unsigned short, int>, subtract_gpu<unsigned short, float>, subtract_gpu<unsigned short, double>},\r
+        {0/*subtract_gpu<short, unsigned char>*/, 0/*subtract_gpu<short, signed char>*/, 0/*subtract_gpu<short, unsigned short>*/, subtract_gpu<short, short>, subtract_gpu<short, int>, subtract_gpu<short, float>, subtract_gpu<short, double>},\r
+        {0/*subtract_gpu<int, unsigned char>*/, 0/*subtract_gpu<int, signed char>*/, 0/*subtract_gpu<int, unsigned short>*/, 0/*subtract_gpu<int, short>*/, subtract_gpu<int, int>, subtract_gpu<int, float>, subtract_gpu<int, double>},\r
+        {0/*subtract_gpu<float, unsigned char>*/, 0/*subtract_gpu<float, signed char>*/, 0/*subtract_gpu<float, unsigned short>*/, 0/*subtract_gpu<float, short>*/, 0/*subtract_gpu<float, int>*/, subtract_gpu<float, float>, subtract_gpu<float, double>},\r
+        {0/*subtract_gpu<double, unsigned char>*/, 0/*subtract_gpu<double, signed char>*/, 0/*subtract_gpu<double, unsigned short>*/, 0/*subtract_gpu<double, short>*/, 0/*subtract_gpu<double, int>*/, 0/*subtract_gpu<double, float>*/, subtract_gpu<double, double>}\r
+    };\r
+\r
+    CV_Assert(src1.type() == src2.type() && src1.size() == src2.size());\r
+    CV_Assert(mask.empty() || (src1.channels() == 1 && mask.size() == src1.size() && mask.type() == CV_8U));\r
+\r
+    if (dtype < 0)\r
+        dtype = src1.depth();\r
+\r
+    dst.create(src1.size(), CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src1.channels()));\r
+\r
+    cudaStream_t stream = StreamAccessor::getStream(s);\r
+\r
+    if (mask.empty() && dst.type() == src1.type() && (src1.depth() == CV_8U || src1.depth() == CV_32S || src1.depth() == CV_32F))\r
+    {\r
+        nppArithmCaller(src2, src1, dst, nppiSub_8u_C1RSfs, nppiSub_8u_C4RSfs, nppiSub_32s_C1R, nppiSub_32f_C1R, stream);\r
+        return;\r
     }\r
-    else\r
-        nppArithmCaller(src1, src2, dst, nppiAdd_8u_C1RSfs, nppiAdd_8u_C4RSfs, nppiAdd_32s_C1R, nppiAdd_32f_C1R, StreamAccessor::getStream(stream));\r
+\r
+    const func_t func = funcs[src1.depth()][dst.depth()];\r
+    CV_Assert(func != 0);\r
+\r
+    func(src1.reshape(1), src2.reshape(1), dst.reshape(1), mask, stream);\r
 }\r
 \r
-namespace cv { namespace gpu { namespace device\r
+void cv::gpu::subtract(const GpuMat& src, const Scalar& sc, GpuMat& dst, const GpuMat& mask, int dtype, Stream& s)\r
 {\r
-    template <typename T>\r
-    void subtractCaller(const DevMem2D src1, const DevMem2D src2, DevMem2D dst, cudaStream_t stream);\r
-}}}\r
+    using namespace cv::gpu::device;\r
 \r
-void cv::gpu::subtract(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)\r
-{\r
-    if (src1.depth() == CV_16S && src2.depth() == CV_16S)\r
+    typedef void (*func_t)(const DevMem2D& src1, double val, const DevMem2D& dst, const PtrStep& mask, cudaStream_t stream);\r
+\r
+    static const func_t funcs[7][7] = \r
     {\r
-        CV_Assert(src1.size() == src2.size());\r
-        dst.create(src1.size(), src1.type());\r
-        device::subtractCaller<short>(src1.reshape(1), src2.reshape(1), dst.reshape(1), StreamAccessor::getStream(stream));\r
+        {subtract_gpu<unsigned char, unsigned char>, 0/*subtract_gpu<unsigned char, signed char>*/, subtract_gpu<unsigned char, unsigned short>, subtract_gpu<unsigned char, short>, subtract_gpu<unsigned char, int>, subtract_gpu<unsigned char, float>, subtract_gpu<unsigned char, double>},\r
+        {0/*subtract_gpu<signed char, unsigned char>*/, 0/*subtract_gpu<signed char, signed char>*/, 0/*subtract_gpu<signed char, unsigned short>*/, 0/*subtract_gpu<signed char, short>*/, 0/*subtract_gpu<signed char, int>*/, 0/*subtract_gpu<signed char, float>*/, 0/*subtract_gpu<signed char, double>*/},\r
+        {0/*subtract_gpu<unsigned short, unsigned char>*/, 0/*subtract_gpu<unsigned short, signed char>*/, subtract_gpu<unsigned short, unsigned short>, 0/*subtract_gpu<unsigned short, short>*/, subtract_gpu<unsigned short, int>, subtract_gpu<unsigned short, float>, subtract_gpu<unsigned short, double>},\r
+        {0/*subtract_gpu<short, unsigned char>*/, 0/*subtract_gpu<short, signed char>*/, 0/*subtract_gpu<short, unsigned short>*/, subtract_gpu<short, short>, subtract_gpu<short, int>, subtract_gpu<short, float>, subtract_gpu<short, double>},\r
+        {0/*subtract_gpu<int, unsigned char>*/, 0/*subtract_gpu<int, signed char>*/, 0/*subtract_gpu<int, unsigned short>*/, 0/*subtract_gpu<int, short>*/, subtract_gpu<int, int>, subtract_gpu<int, float>, subtract_gpu<int, double>},\r
+        {0/*subtract_gpu<float, unsigned char>*/, 0/*subtract_gpu<float, signed char>*/, 0/*subtract_gpu<float, unsigned short>*/, 0/*subtract_gpu<float, short>*/, 0/*subtract_gpu<float, int>*/, subtract_gpu<float, float>, subtract_gpu<float, double>},\r
+        {0/*subtract_gpu<double, unsigned char>*/, 0/*subtract_gpu<double, signed char>*/, 0/*subtract_gpu<double, unsigned short>*/, 0/*subtract_gpu<double, short>*/, 0/*subtract_gpu<double, int>*/, 0/*subtract_gpu<double, float>*/, subtract_gpu<double, double>}\r
+    };\r
+\r
+    CV_Assert(src.channels() == 1 || src.type() == CV_32FC2);\r
+    CV_Assert(mask.empty() || (src.channels() == 1 && mask.size() == src.size() && mask.type() == CV_8U));\r
+\r
+    if (dtype < 0)\r
+        dtype = src.depth();\r
+\r
+    dst.create(src.size(), CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src.channels()));\r
+\r
+    cudaStream_t stream = StreamAccessor::getStream(s);\r
+\r
+    if (mask.empty() && dst.type() == src.type() && src.depth() == CV_32F)\r
+    {\r
+        NppiSize sz;\r
+        sz.width  = src.cols;\r
+        sz.height = src.rows;\r
+\r
+        NppStreamHandler h(stream);\r
+\r
+        if (src.type() == CV_32FC1)\r
+        {\r
+            nppSafeCall( nppiSubC_32f_C1R(src.ptr<Npp32f>(), static_cast<int>(src.step), static_cast<Npp32f>(sc.val[0]), \r
+                dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz) );\r
+        }\r
+        else\r
+        {\r
+            Npp32fc nValue;\r
+            nValue.re = static_cast<Npp32f>(sc.val[0]);\r
+            nValue.im = static_cast<Npp32f>(sc.val[1]);\r
+            nppSafeCall( nppiSubC_32fc_C1R(src.ptr<Npp32fc>(), static_cast<int>(src.step), nValue, \r
+                dst.ptr<Npp32fc>(), static_cast<int>(dst.step), sz) );\r
+        }\r
+\r
+        if (stream == 0)\r
+            cudaSafeCall( cudaDeviceSynchronize() );\r
+\r
+        return;\r
     }\r
-    else\r
-        nppArithmCaller(src2, src1, dst, nppiSub_8u_C1RSfs, nppiSub_8u_C4RSfs, nppiSub_32s_C1R, nppiSub_32f_C1R, StreamAccessor::getStream(stream));\r
+\r
+    const func_t func = funcs[src.depth()][dst.depth()];\r
+    CV_Assert(func != 0);\r
+\r
+    func(src, sc.val[0], dst, mask, stream);\r
 }\r
 \r
+////////////////////////////////////////////////////////////////////////\r
+// multiply\r
+\r
 namespace cv { namespace gpu { namespace device\r
 {\r
     void multiply_gpu(const DevMem2D_<uchar4>& src1, const DevMem2Df& src2, const DevMem2D_<uchar4>& dst, cudaStream_t stream);\r
     void multiply_gpu(const DevMem2D_<short4>& src1, const DevMem2Df& src2, const DevMem2D_<short4>& dst, cudaStream_t stream);\r
 \r
-    template <typename T, typename D>\r
-    void multiplyScalar_gpu(const DevMem2D& src, float scale, const DevMem2D& dst, cudaStream_t stream);\r
+    template <typename T, typename D> \r
+    void multiply_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    template <typename T, typename D> \r
+    void multiply_gpu(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
 }}}\r
 \r
-void cv::gpu::multiply(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)\r
+void cv::gpu::multiply(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, double scale, int dtype, Stream& s)\r
 {\r
+    using namespace cv::gpu::device;\r
+\r
+    typedef void (*func_t)(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    static const func_t funcs[7][7] = \r
+    {\r
+        {multiply_gpu<unsigned char, unsigned char>, 0/*multiply_gpu<unsigned char, signed char>*/, multiply_gpu<unsigned char, unsigned short>, multiply_gpu<unsigned char, short>, multiply_gpu<unsigned char, int>, multiply_gpu<unsigned char, float>, multiply_gpu<unsigned char, double>},\r
+        {0/*multiply_gpu<signed char, unsigned char>*/, 0/*multiply_gpu<signed char, signed char>*/, 0/*multiply_gpu<signed char, unsigned short>*/, 0/*multiply_gpu<signed char, short>*/, 0/*multiply_gpu<signed char, int>*/, 0/*multiply_gpu<signed char, float>*/, 0/*multiply_gpu<signed char, double>*/},\r
+        {0/*multiply_gpu<unsigned short, unsigned char>*/, 0/*multiply_gpu<unsigned short, signed char>*/, multiply_gpu<unsigned short, unsigned short>, 0/*multiply_gpu<unsigned short, short>*/, multiply_gpu<unsigned short, int>, multiply_gpu<unsigned short, float>, multiply_gpu<unsigned short, double>},\r
+        {0/*multiply_gpu<short, unsigned char>*/, 0/*multiply_gpu<short, signed char>*/, 0/*multiply_gpu<short, unsigned short>*/, multiply_gpu<short, short>, multiply_gpu<short, int>, multiply_gpu<short, float>, multiply_gpu<short, double>},\r
+        {0/*multiply_gpu<int, unsigned char>*/, 0/*multiply_gpu<int, signed char>*/, 0/*multiply_gpu<int, unsigned short>*/, 0/*multiply_gpu<int, short>*/, multiply_gpu<int, int>, multiply_gpu<int, float>, multiply_gpu<int, double>},\r
+        {0/*multiply_gpu<float, unsigned char>*/, 0/*multiply_gpu<float, signed char>*/, 0/*multiply_gpu<float, unsigned short>*/, 0/*multiply_gpu<float, short>*/, 0/*multiply_gpu<float, int>*/, multiply_gpu<float, float>, multiply_gpu<float, double>},\r
+        {0/*multiply_gpu<double, unsigned char>*/, 0/*multiply_gpu<double, signed char>*/, 0/*multiply_gpu<double, unsigned short>*/, 0/*multiply_gpu<double, short>*/, 0/*multiply_gpu<double, int>*/, 0/*multiply_gpu<double, float>*/, multiply_gpu<double, double>}\r
+    };\r
+\r
+    cudaStream_t stream = StreamAccessor::getStream(s);\r
+\r
     if (src1.type() == CV_8UC4 && src2.type() == CV_32FC1)\r
     {\r
         CV_Assert(src1.size() == src2.size());\r
 \r
         dst.create(src1.size(), src1.type());\r
 \r
-        device::multiply_gpu(static_cast<DevMem2D_<uchar4> >(src1), static_cast<DevMem2Df>(src2), \r
-                             static_cast<DevMem2D_<uchar4> >(dst), StreamAccessor::getStream(stream));\r
+        multiply_gpu(static_cast<DevMem2D_<uchar4> >(src1), static_cast<DevMem2Df>(src2), static_cast<DevMem2D_<uchar4> >(dst), stream);\r
     }\r
     else if (src1.type() == CV_16SC4 && src2.type() == CV_32FC1)\r
     {\r
@@ -237,60 +392,124 @@ void cv::gpu::multiply(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stre
 \r
         dst.create(src1.size(), src1.type());\r
 \r
-        device::multiply_gpu(static_cast<DevMem2D_<short4> >(src1), static_cast<DevMem2Df>(src2), \r
-                             static_cast<DevMem2D_<short4> >(dst), StreamAccessor::getStream(stream));\r
+        multiply_gpu(static_cast<DevMem2D_<short4> >(src1), static_cast<DevMem2Df>(src2), static_cast<DevMem2D_<short4> >(dst), stream);\r
     }\r
     else\r
-        nppArithmCaller(src1, src2, dst, nppiMul_8u_C1RSfs, nppiMul_8u_C4RSfs, nppiMul_32s_C1R, nppiMul_32f_C1R, StreamAccessor::getStream(stream));\r
+    {\r
+        CV_Assert(src1.type() == src2.type() && src1.size() == src2.size());\r
+\r
+        if (dtype < 0)\r
+            dtype = src1.depth();\r
+\r
+        dst.create(src1.size(), CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src1.channels()));\r
+\r
+\r
+        if (scale == 1 && dst.type() == src1.type() && (src1.depth() == CV_8U || src1.depth() == CV_32S || src1.depth() == CV_32F))\r
+        {\r
+            nppArithmCaller(src2, src1, dst, nppiMul_8u_C1RSfs, nppiMul_8u_C4RSfs, nppiMul_32s_C1R, nppiMul_32f_C1R, stream);\r
+            return;\r
+        }\r
+\r
+        const func_t func = funcs[src1.depth()][dst.depth()];\r
+        CV_Assert(func != 0);\r
+\r
+        func(src1.reshape(1), src2.reshape(1), dst.reshape(1), scale, stream);\r
+    }\r
 }\r
 \r
-void cv::gpu::multiply(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream& stream)\r
+void cv::gpu::multiply(const GpuMat& src, const Scalar& sc, GpuMat& dst, double scale, int dtype, Stream& s)\r
 {\r
-    if (src.depth() == CV_8U)\r
-    {\r
-        dst.create(src.size(), src.type());\r
+    using namespace cv::gpu::device;\r
 \r
-        device::multiplyScalar_gpu<unsigned char, unsigned char>(src.reshape(1), (float)(sc[0]), dst, StreamAccessor::getStream(stream));\r
-    }\r
-    else\r
+    typedef void (*func_t)(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    static const func_t funcs[7][7] = \r
     {\r
-        CV_Assert(src.type() == CV_32FC1);\r
+        {multiply_gpu<unsigned char, unsigned char>, 0/*multiply_gpu<unsigned char, signed char>*/, multiply_gpu<unsigned char, unsigned short>, multiply_gpu<unsigned char, short>, multiply_gpu<unsigned char, int>, multiply_gpu<unsigned char, float>, multiply_gpu<unsigned char, double>},\r
+        {0/*multiply_gpu<signed char, unsigned char>*/, 0/*multiply_gpu<signed char, signed char>*/, 0/*multiply_gpu<signed char, unsigned short>*/, 0/*multiply_gpu<signed char, short>*/, 0/*multiply_gpu<signed char, int>*/, 0/*multiply_gpu<signed char, float>*/, 0/*multiply_gpu<signed char, double>*/},\r
+        {0/*multiply_gpu<unsigned short, unsigned char>*/, 0/*multiply_gpu<unsigned short, signed char>*/, multiply_gpu<unsigned short, unsigned short>, 0/*multiply_gpu<unsigned short, short>*/, multiply_gpu<unsigned short, int>, multiply_gpu<unsigned short, float>, multiply_gpu<unsigned short, double>},\r
+        {0/*multiply_gpu<short, unsigned char>*/, 0/*multiply_gpu<short, signed char>*/, 0/*multiply_gpu<short, unsigned short>*/, multiply_gpu<short, short>, multiply_gpu<short, int>, multiply_gpu<short, float>, multiply_gpu<short, double>},\r
+        {0/*multiply_gpu<int, unsigned char>*/, 0/*multiply_gpu<int, signed char>*/, 0/*multiply_gpu<int, unsigned short>*/, 0/*multiply_gpu<int, short>*/, multiply_gpu<int, int>, multiply_gpu<int, float>, multiply_gpu<int, double>},\r
+        {0/*multiply_gpu<float, unsigned char>*/, 0/*multiply_gpu<float, signed char>*/, 0/*multiply_gpu<float, unsigned short>*/, 0/*multiply_gpu<float, short>*/, 0/*multiply_gpu<float, int>*/, multiply_gpu<float, float>, multiply_gpu<float, double>},\r
+        {0/*multiply_gpu<double, unsigned char>*/, 0/*multiply_gpu<double, signed char>*/, 0/*multiply_gpu<double, unsigned short>*/, 0/*multiply_gpu<double, short>*/, 0/*multiply_gpu<double, int>*/, 0/*multiply_gpu<double, float>*/, multiply_gpu<double, double>}\r
+    };\r
 \r
-        dst.create(src.size(), src.type());\r
+    CV_Assert(src.channels() == 1);\r
+\r
+    if (dtype < 0)\r
+        dtype = src.depth();\r
+\r
+    dst.create(src.size(), CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src.channels()));\r
+\r
+    cudaStream_t stream = StreamAccessor::getStream(s);\r
 \r
+    if (dst.type() == src.type() && src.type() == CV_32FC1 && scale == 1)\r
+    {\r
         NppiSize sz;\r
         sz.width  = src.cols;\r
         sz.height = src.rows;\r
 \r
-        cudaStream_t cudaStream = StreamAccessor::getStream(stream);\r
-\r
-        NppStreamHandler h(cudaStream);\r
+        NppStreamHandler h(stream);\r
 \r
-        nppSafeCall( nppiMulC_32f_C1R(src.ptr<Npp32f>(), static_cast<int>(src.step), (Npp32f)sc[0], dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz) );\r
+        nppSafeCall( nppiMulC_32f_C1R(src.ptr<Npp32f>(), static_cast<int>(src.step), static_cast<Npp32f>(sc.val[0]), \r
+            dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz) );\r
 \r
-        if (cudaStream == 0)\r
+        if (stream == 0)\r
             cudaSafeCall( cudaDeviceSynchronize() );\r
+\r
+        return;\r
     }\r
+\r
+    const func_t func = funcs[src.depth()][dst.depth()];\r
+    CV_Assert(func != 0);\r
+\r
+    func(src, sc.val[0], dst, scale, stream);\r
 }\r
 \r
+////////////////////////////////////////////////////////////////////////\r
+// divide\r
 \r
 namespace cv { namespace gpu { namespace device\r
 {\r
     void divide_gpu(const DevMem2D_<uchar4>& src1, const DevMem2Df& src2, const DevMem2D_<uchar4>& dst, cudaStream_t stream);\r
     void divide_gpu(const DevMem2D_<short4>& src1, const DevMem2Df& src2, const DevMem2D_<short4>& dst, cudaStream_t stream);\r
+    \r
+    template <typename T, typename D> \r
+    void divide_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    \r
+    template <typename T, typename D> \r
+    void divide_gpu(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+    \r
+    template <typename T, typename D> \r
+    void divide_gpu(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
 }}}\r
 \r
-\r
-void cv::gpu::divide(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& stream)\r
+void cv::gpu::divide(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, double scale, int dtype, Stream& s)\r
 {\r
+    using namespace cv::gpu::device;\r
+\r
+    typedef void (*func_t)(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    static const func_t funcs[7][7] = \r
+    {\r
+        {divide_gpu<unsigned char, unsigned char>, 0/*divide_gpu<unsigned char, signed char>*/, divide_gpu<unsigned char, unsigned short>, divide_gpu<unsigned char, short>, divide_gpu<unsigned char, int>, divide_gpu<unsigned char, float>, divide_gpu<unsigned char, double>},\r
+        {0/*divide_gpu<signed char, unsigned char>*/, 0/*divide_gpu<signed char, signed char>*/, 0/*divide_gpu<signed char, unsigned short>*/, 0/*divide_gpu<signed char, short>*/, 0/*divide_gpu<signed char, int>*/, 0/*divide_gpu<signed char, float>*/, 0/*divide_gpu<signed char, double>*/},\r
+        {0/*divide_gpu<unsigned short, unsigned char>*/, 0/*divide_gpu<unsigned short, signed char>*/, divide_gpu<unsigned short, unsigned short>, 0/*divide_gpu<unsigned short, short>*/, divide_gpu<unsigned short, int>, divide_gpu<unsigned short, float>, divide_gpu<unsigned short, double>},\r
+        {0/*divide_gpu<short, unsigned char>*/, 0/*divide_gpu<short, signed char>*/, 0/*divide_gpu<short, unsigned short>*/, divide_gpu<short, short>, divide_gpu<short, int>, divide_gpu<short, float>, divide_gpu<short, double>},\r
+        {0/*divide_gpu<int, unsigned char>*/, 0/*divide_gpu<int, signed char>*/, 0/*divide_gpu<int, unsigned short>*/, 0/*divide_gpu<int, short>*/, divide_gpu<int, int>, divide_gpu<int, float>, divide_gpu<int, double>},\r
+        {0/*divide_gpu<float, unsigned char>*/, 0/*divide_gpu<float, signed char>*/, 0/*divide_gpu<float, unsigned short>*/, 0/*divide_gpu<float, short>*/, 0/*divide_gpu<float, int>*/, divide_gpu<float, float>, divide_gpu<float, double>},\r
+        {0/*divide_gpu<double, unsigned char>*/, 0/*divide_gpu<double, signed char>*/, 0/*divide_gpu<double, unsigned short>*/, 0/*divide_gpu<double, short>*/, 0/*divide_gpu<double, int>*/, 0/*divide_gpu<double, float>*/, divide_gpu<double, double>}\r
+    };\r
+\r
+    cudaStream_t stream = StreamAccessor::getStream(s);\r
+\r
     if (src1.type() == CV_8UC4 && src2.type() == CV_32FC1)\r
     {\r
         CV_Assert(src1.size() == src2.size());\r
 \r
         dst.create(src1.size(), src1.type());\r
 \r
-        device::divide_gpu(static_cast<DevMem2D_<uchar4> >(src1), static_cast<DevMem2Df>(src2),\r
-                           static_cast<DevMem2D_<uchar4> >(dst), StreamAccessor::getStream(stream));\r
+        multiply_gpu(static_cast<DevMem2D_<uchar4> >(src1), static_cast<DevMem2Df>(src2), static_cast<DevMem2D_<uchar4> >(dst), stream);\r
     }\r
     else if (src1.type() == CV_16SC4 && src2.type() == CV_32FC1)\r
     {\r
@@ -298,117 +517,236 @@ void cv::gpu::divide(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream
 \r
         dst.create(src1.size(), src1.type());\r
 \r
-        device::divide_gpu(static_cast<DevMem2D_<short4> >(src1), static_cast<DevMem2Df>(src2),\r
-                           static_cast<DevMem2D_<short4> >(dst), StreamAccessor::getStream(stream));\r
+        multiply_gpu(static_cast<DevMem2D_<short4> >(src1), static_cast<DevMem2Df>(src2), static_cast<DevMem2D_<short4> >(dst), stream);\r
     }\r
     else\r
-        nppArithmCaller(src2, src1, dst, nppiDiv_8u_C1RSfs, nppiDiv_8u_C4RSfs, nppiDiv_32s_C1R, nppiDiv_32f_C1R, StreamAccessor::getStream(stream));\r
-}\r
+    {\r
+        CV_Assert(src1.type() == src2.type() && src1.size() == src2.size());\r
 \r
-void cv::gpu::add(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream& stream)\r
-{\r
-    typedef void (*caller_t)(const GpuMat& src, const Scalar& sc, GpuMat& dst, cudaStream_t stream);\r
-    static const caller_t callers[] = {0, NppArithmScalar<1, nppiAddC_32f_C1R>::calc, NppArithmScalar<2, nppiAddC_32fc_C1R>::calc};\r
+        if (dtype < 0)\r
+            dtype = src1.depth();\r
+\r
+        dst.create(src1.size(), CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src1.channels()));\r
+\r
+\r
+        if (scale == 1 && dst.type() == src1.type() && (src1.depth() == CV_8U || src1.depth() == CV_32S || src1.depth() == CV_32F))\r
+        {\r
+            nppArithmCaller(src2, src1, dst, nppiDiv_8u_C1RSfs, nppiDiv_8u_C4RSfs, nppiDiv_32s_C1R, nppiDiv_32f_C1R, stream);\r
+            return;\r
+        }\r
 \r
-    CV_Assert(src.type() == CV_32FC1 || src.type() == CV_32FC2);\r
+        const func_t func = funcs[src1.depth()][dst.depth()];\r
+        CV_Assert(func != 0);\r
 \r
-    callers[src.channels()](src, sc, dst, StreamAccessor::getStream(stream));\r
+        func(src1.reshape(1), src2.reshape(1), dst.reshape(1), scale, stream);\r
+    }\r
 }\r
 \r
-void cv::gpu::subtract(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream& stream)\r
+void cv::gpu::divide(const GpuMat& src, const Scalar& sc, GpuMat& dst, double scale, int dtype, Stream& s)\r
 {\r
-    typedef void (*caller_t)(const GpuMat& src, const Scalar& sc, GpuMat& dst, cudaStream_t stream);\r
-    static const caller_t callers[] = {0, NppArithmScalar<1, nppiSubC_32f_C1R>::calc, NppArithmScalar<2, nppiSubC_32fc_C1R>::calc};\r
+    using namespace cv::gpu::device;\r
+\r
+    typedef void (*func_t)(const DevMem2D& src1, double val, const DevMem2D& dst, double scale, cudaStream_t stream);\r
+\r
+    static const func_t funcs[7][7] = \r
+    {\r
+        {divide_gpu<unsigned char, unsigned char>, 0/*divide_gpu<unsigned char, signed char>*/, divide_gpu<unsigned char, unsigned short>, divide_gpu<unsigned char, short>, divide_gpu<unsigned char, int>, divide_gpu<unsigned char, float>, divide_gpu<unsigned char, double>},\r
+        {0/*divide_gpu<signed char, unsigned char>*/, 0/*divide_gpu<signed char, signed char>*/, 0/*divide_gpu<signed char, unsigned short>*/, 0/*divide_gpu<signed char, short>*/, 0/*divide_gpu<signed char, int>*/, 0/*divide_gpu<signed char, float>*/, 0/*divide_gpu<signed char, double>*/},\r
+        {0/*divide_gpu<unsigned short, unsigned char>*/, 0/*divide_gpu<unsigned short, signed char>*/, divide_gpu<unsigned short, unsigned short>, 0/*divide_gpu<unsigned short, short>*/, divide_gpu<unsigned short, int>, divide_gpu<unsigned short, float>, divide_gpu<unsigned short, double>},\r
+        {0/*divide_gpu<short, unsigned char>*/, 0/*divide_gpu<short, signed char>*/, 0/*divide_gpu<short, unsigned short>*/, divide_gpu<short, short>, divide_gpu<short, int>, divide_gpu<short, float>, divide_gpu<short, double>},\r
+        {0/*divide_gpu<int, unsigned char>*/, 0/*divide_gpu<int, signed char>*/, 0/*divide_gpu<int, unsigned short>*/, 0/*divide_gpu<int, short>*/, divide_gpu<int, int>, divide_gpu<int, float>, divide_gpu<int, double>},\r
+        {0/*divide_gpu<float, unsigned char>*/, 0/*divide_gpu<float, signed char>*/, 0/*divide_gpu<float, unsigned short>*/, 0/*divide_gpu<float, short>*/, 0/*divide_gpu<float, int>*/, divide_gpu<float, float>, divide_gpu<float, double>},\r
+        {0/*divide_gpu<double, unsigned char>*/, 0/*divide_gpu<double, signed char>*/, 0/*divide_gpu<double, unsigned short>*/, 0/*divide_gpu<double, short>*/, 0/*divide_gpu<double, int>*/, 0/*divide_gpu<double, float>*/, divide_gpu<double, double>}\r
+    };\r
+\r
+    CV_Assert(src.channels() == 1);\r
+\r
+    if (dtype < 0)\r
+        dtype = src.depth();\r
+\r
+    dst.create(src.size(), CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src.channels()));\r
+\r
+    cudaStream_t stream = StreamAccessor::getStream(s);\r
+\r
+    if (dst.type() == src.type() && src.type() == CV_32FC1 && scale == 1)\r
+    {\r
+        NppiSize sz;\r
+        sz.width  = src.cols;\r
+        sz.height = src.rows;\r
+\r
+        NppStreamHandler h(stream);\r
+\r
+        nppSafeCall( nppiDivC_32f_C1R(src.ptr<Npp32f>(), static_cast<int>(src.step), static_cast<Npp32f>(sc.val[0]), \r
+            dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz) );\r
 \r
-    CV_Assert(src.type() == CV_32FC1 || src.type() == CV_32FC2);\r
+        if (stream == 0)\r
+            cudaSafeCall( cudaDeviceSynchronize() );\r
+\r
+        return;\r
+    }\r
+\r
+    const func_t func = funcs[src.depth()][dst.depth()];\r
+    CV_Assert(func != 0);\r
 \r
-    callers[src.channels()](src, sc, dst, StreamAccessor::getStream(stream));\r
+    func(src, sc.val[0], dst, scale, stream);\r
 }\r
 \r
-void cv::gpu::divide(const GpuMat& src, const Scalar& sc, GpuMat& dst, Stream& stream)\r
+void cv::gpu::divide(double scale, const GpuMat& src, GpuMat& dst, int dtype, Stream& s)\r
 {\r
-    CV_Assert(src.type() == CV_32FC1);\r
+    using namespace cv::gpu::device;\r
 \r
-    dst.create(src.size(), src.type());\r
+    typedef void (*func_t)(double scalar, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
 \r
-    NppiSize sz;\r
-    sz.width  = src.cols;\r
-    sz.height = src.rows;\r
+    static const func_t funcs[7][7] = \r
+    {\r
+        {divide_gpu<unsigned char, unsigned char>, 0/*divide_gpu<unsigned char, signed char>*/, divide_gpu<unsigned char, unsigned short>, divide_gpu<unsigned char, short>, divide_gpu<unsigned char, int>, divide_gpu<unsigned char, float>, divide_gpu<unsigned char, double>},\r
+        {0/*divide_gpu<signed char, unsigned char>*/, 0/*divide_gpu<signed char, signed char>*/, 0/*divide_gpu<signed char, unsigned short>*/, 0/*divide_gpu<signed char, short>*/, 0/*divide_gpu<signed char, int>*/, 0/*divide_gpu<signed char, float>*/, 0/*divide_gpu<signed char, double>*/},\r
+        {0/*divide_gpu<unsigned short, unsigned char>*/, 0/*divide_gpu<unsigned short, signed char>*/, divide_gpu<unsigned short, unsigned short>, 0/*divide_gpu<unsigned short, short>*/, divide_gpu<unsigned short, int>, divide_gpu<unsigned short, float>, divide_gpu<unsigned short, double>},\r
+        {0/*divide_gpu<short, unsigned char>*/, 0/*divide_gpu<short, signed char>*/, 0/*divide_gpu<short, unsigned short>*/, divide_gpu<short, short>, divide_gpu<short, int>, divide_gpu<short, float>, divide_gpu<short, double>},\r
+        {0/*divide_gpu<int, unsigned char>*/, 0/*divide_gpu<int, signed char>*/, 0/*divide_gpu<int, unsigned short>*/, 0/*divide_gpu<int, short>*/, divide_gpu<int, int>, divide_gpu<int, float>, divide_gpu<int, double>},\r
+        {0/*divide_gpu<float, unsigned char>*/, 0/*divide_gpu<float, signed char>*/, 0/*divide_gpu<float, unsigned short>*/, 0/*divide_gpu<float, short>*/, 0/*divide_gpu<float, int>*/, divide_gpu<float, float>, divide_gpu<float, double>},\r
+        {0/*divide_gpu<double, unsigned char>*/, 0/*divide_gpu<double, signed char>*/, 0/*divide_gpu<double, unsigned short>*/, 0/*divide_gpu<double, short>*/, 0/*divide_gpu<double, int>*/, 0/*divide_gpu<double, float>*/, divide_gpu<double, double>}\r
+    };\r
 \r
-    cudaStream_t cudaStream = StreamAccessor::getStream(stream);\r
+    CV_Assert(src.channels() == 1);\r
 \r
-    NppStreamHandler h(cudaStream);\r
+    if (dtype < 0)\r
+        dtype = src.depth();\r
 \r
-    nppSafeCall( nppiDivC_32f_C1R(src.ptr<Npp32f>(), static_cast<int>(src.step), (Npp32f)sc[0], dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz) );\r
+    dst.create(src.size(), CV_MAKE_TYPE(CV_MAT_DEPTH(dtype), src.channels()));\r
 \r
-    if (cudaStream == 0)\r
-        cudaSafeCall( cudaDeviceSynchronize() );\r
-}\r
+    cudaStream_t stream = StreamAccessor::getStream(s);\r
 \r
+    const func_t func = funcs[src.depth()][dst.depth()];\r
+    CV_Assert(func != 0);\r
+\r
+    func(scale, src, dst, stream);\r
+}\r
 \r
 //////////////////////////////////////////////////////////////////////////////\r
-// Absolute difference\r
+// absdiff\r
+\r
+namespace cv { namespace gpu { namespace device\r
+{\r
+    template <typename T> \r
+    void absdiff_gpu(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    \r
+    template <typename T> \r
+    void absdiff_gpu(const DevMem2D& src1, double val, const DevMem2D& dst, cudaStream_t stream);\r
+}}}\r
 \r
 void cv::gpu::absdiff(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, Stream& s)\r
 {\r
-    CV_DbgAssert(src1.size() == src2.size() && src1.type() == src2.type());\r
+    using namespace cv::gpu::device;\r
 \r
-    CV_Assert(src1.type() == CV_8UC1 || src1.type() == CV_8UC4 || src1.type() == CV_32SC1 || src1.type() == CV_32FC1);\r
+    typedef void (*func_t)(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+\r
+    static const func_t funcs[] = \r
+    {\r
+        0/*absdiff_gpu<unsigned char>*/, absdiff_gpu<signed char>, absdiff_gpu<unsigned short>, absdiff_gpu<short>, 0/*absdiff_gpu<int>*/, 0/*absdiff_gpu<float>*/, absdiff_gpu<double>\r
+    };\r
+\r
+    CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
 \r
     dst.create( src1.size(), src1.type() );\r
 \r
+    cudaStream_t stream = StreamAccessor::getStream(s);\r
+\r
     NppiSize sz;\r
-    sz.width  = src1.cols;\r
+    sz.width  = src1.cols * src1.channels();\r
     sz.height = src1.rows;\r
 \r
-    cudaStream_t stream = StreamAccessor::getStream(s);\r
+    if (src1.depth() == CV_8U && (src1.cols * src1.channels()) % 4 == 0)\r
+    {\r
+        NppStreamHandler h(stream);\r
 \r
-    NppStreamHandler h(stream);\r
+        sz.width /= 4;\r
 \r
-    switch (src1.type())\r
+        nppSafeCall( nppiAbsDiff_8u_C4R(src1.ptr<Npp8u>(), static_cast<int>(src1.step), src2.ptr<Npp8u>(), static_cast<int>(src2.step), \r
+            dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz) );\r
+\r
+        if (stream == 0)\r
+            cudaSafeCall( cudaDeviceSynchronize() );\r
+    }\r
+    else if (src1.depth() == CV_8U)\r
     {\r
-    case CV_8UC1:\r
+        NppStreamHandler h(stream);\r
+\r
         nppSafeCall( nppiAbsDiff_8u_C1R(src1.ptr<Npp8u>(), static_cast<int>(src1.step), src2.ptr<Npp8u>(), static_cast<int>(src2.step), \r
             dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz) );\r
-        break;\r
-    case CV_8UC4:\r
-        nppSafeCall( nppiAbsDiff_8u_C4R(src1.ptr<Npp8u>(), static_cast<int>(src1.step), src2.ptr<Npp8u>(), static_cast<int>(src2.step), \r
-            dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz) );\r
-        break;\r
-    case CV_32SC1:\r
+\r
+        if (stream == 0)\r
+            cudaSafeCall( cudaDeviceSynchronize() );\r
+    }\r
+    else if (src1.depth() == CV_32S)\r
+    {\r
+        NppStreamHandler h(stream);\r
+\r
         nppSafeCall( nppiAbsDiff_32s_C1R(src1.ptr<Npp32s>(), static_cast<int>(src1.step), src2.ptr<Npp32s>(), static_cast<int>(src2.step), \r
             dst.ptr<Npp32s>(), static_cast<int>(dst.step), sz) );\r
-        break;\r
-    case CV_32FC1:\r
+\r
+        if (stream == 0)\r
+            cudaSafeCall( cudaDeviceSynchronize() );\r
+    }\r
+    else if (src1.depth() == CV_32F)\r
+    {\r
+        NppStreamHandler h(stream);\r
+\r
         nppSafeCall( nppiAbsDiff_32f_C1R(src1.ptr<Npp32f>(), static_cast<int>(src1.step), src2.ptr<Npp32f>(), static_cast<int>(src2.step), \r
             dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz) );\r
-        break;\r
-    default:\r
-        CV_Assert(!"Unsupported source type");\r
+\r
+        if (stream == 0)\r
+            cudaSafeCall( cudaDeviceSynchronize() );\r
     }\r
+    else\r
+    {\r
+        const func_t func = funcs[src1.depth()];\r
+        CV_Assert(func != 0);\r
 \r
-    if (stream == 0)\r
-        cudaSafeCall( cudaDeviceSynchronize() );\r
+        func(src1.reshape(1), src2.reshape(1), dst.reshape(1), stream);\r
+    }\r
 }\r
 \r
 void cv::gpu::absdiff(const GpuMat& src1, const Scalar& src2, GpuMat& dst, Stream& s)\r
 {\r
-    CV_Assert(src1.type() == CV_32FC1);\r
+    using namespace cv::gpu::device;\r
 \r
-    dst.create( src1.size(), src1.type() );\r
+    typedef void (*func_t)(const DevMem2D& src1, double val, const DevMem2D& dst, cudaStream_t stream);\r
 \r
-    NppiSize sz;\r
-    sz.width  = src1.cols;\r
-    sz.height = src1.rows;\r
+    static const func_t funcs[] = \r
+    {\r
+        absdiff_gpu<unsigned char>, absdiff_gpu<signed char>, absdiff_gpu<unsigned short>, absdiff_gpu<short>,absdiff_gpu<int>, 0/*absdiff_gpu<float>*/, absdiff_gpu<double>\r
+    };\r
+\r
+    CV_Assert(src1.channels() == 1);\r
+\r
+    dst.create(src1.size(), src1.type());\r
 \r
     cudaStream_t stream = StreamAccessor::getStream(s);\r
 \r
-    NppStreamHandler h(stream);\r
+    if (src1.type() == CV_32FC1)\r
+    {\r
+        NppiSize sz;\r
+        sz.width  = src1.cols;\r
+        sz.height = src1.rows;\r
+\r
+        cudaStream_t stream = StreamAccessor::getStream(s);\r
 \r
-    nppSafeCall( nppiAbsDiffC_32f_C1R(src1.ptr<Npp32f>(), static_cast<int>(src1.step), dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz, (Npp32f)src2[0]) );\r
+        NppStreamHandler h(stream);\r
+\r
+        nppSafeCall( nppiAbsDiffC_32f_C1R(src1.ptr<Npp32f>(), static_cast<int>(src1.step), \r
+            dst.ptr<Npp32f>(), static_cast<int>(dst.step), sz, static_cast<Npp32f>(src2.val[0])) );\r
+\r
+        if (stream == 0)\r
+            cudaSafeCall( cudaDeviceSynchronize() );\r
 \r
-    if (stream == 0)\r
-        cudaSafeCall( cudaDeviceSynchronize() );\r
+        return;\r
+    }\r
+\r
+    const func_t func = funcs[src1.depth()];\r
+    CV_Assert(func != 0);\r
+\r
+    func(src1, src2.val[0], dst, stream);\r
 }\r
 \r
 \r
@@ -417,62 +755,74 @@ void cv::gpu::absdiff(const GpuMat& src1, const Scalar& src2, GpuMat& dst, Strea
 \r
 namespace cv { namespace gpu { namespace device\r
 {\r
-    void compare_ne_8uc4(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
-    void compare_ne_32f(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template <typename T> void compare_eq(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template <typename T> void compare_ne(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template <typename T> void compare_lt(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
+    template <typename T> void compare_le(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
 }}}\r
 \r
-void cv::gpu::compare(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, int cmpop, Stream& s)\r
+void cv::gpu::compare(const GpuMat& src1, const GpuMat& src2, GpuMat& dst, int cmpop, Stream& stream)\r
 {\r
-    CV_DbgAssert(src1.size() == src2.size() && src1.type() == src2.type());\r
-\r
-    CV_Assert(src1.type() == CV_8UC4 || src1.type() == CV_32FC1);\r
-\r
-    dst.create( src1.size(), CV_8UC1 );\r
-\r
-    static const NppCmpOp nppCmpOp[] = { NPP_CMP_EQ, NPP_CMP_GREATER, NPP_CMP_GREATER_EQ, NPP_CMP_LESS, NPP_CMP_LESS_EQ };\r
-\r
-    NppiSize sz;\r
-    sz.width  = src1.cols;\r
-    sz.height = src1.rows;\r
+    using namespace cv::gpu::device;\r
 \r
-    cudaStream_t stream = StreamAccessor::getStream(s);\r
+    typedef void (*func_t)(const DevMem2D& src1, const DevMem2D& src2, const DevMem2D& dst, cudaStream_t stream);\r
 \r
-    if (src1.type() == CV_8UC4)\r
+    static const func_t funcs[7][4] = \r
     {\r
-        if (cmpop != CMP_NE)\r
-        {\r
-            NppStreamHandler h(stream);\r
+        {compare_eq<unsigned char>, compare_ne<unsigned char>, compare_lt<unsigned char>, compare_le<unsigned char>},\r
+        {compare_eq<signed char>, compare_ne<signed char>, compare_lt<signed char>, compare_le<signed char>},\r
+        {compare_eq<unsigned short>, compare_ne<unsigned short>, compare_lt<unsigned short>, compare_le<unsigned short>},\r
+        {compare_eq<short>, compare_ne<short>, compare_lt<short>, compare_le<short>},\r
+        {compare_eq<int>, compare_ne<int>, compare_lt<int>, compare_le<int>},\r
+        {compare_eq<float>, compare_ne<float>, compare_lt<float>, compare_le<float>},\r
+        {compare_eq<double>, compare_ne<double>, compare_lt<double>, compare_le<double>}\r
+    };\r
 \r
-            nppSafeCall( nppiCompare_8u_C4R(src1.ptr<Npp8u>(), static_cast<int>(src1.step),\r
-                src2.ptr<Npp8u>(), static_cast<int>(src2.step),\r
-                dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz, nppCmpOp[cmpop]) );\r
+    CV_Assert(src1.size() == src2.size() && src1.type() == src2.type());\r
 \r
-            if (stream == 0)\r
-                cudaSafeCall( cudaDeviceSynchronize() );\r
-        }\r
-        else\r
-        {\r
-            device::compare_ne_8uc4(src1, src2, dst, stream);\r
-        }\r
-    }\r
-    else\r
+    int code;\r
+    const GpuMat* psrc1;\r
+    const GpuMat* psrc2;\r
+\r
+    switch (cmpop)\r
     {\r
-        if (cmpop != CMP_NE)\r
-        {\r
-            NppStreamHandler h(stream);\r
+    case CMP_EQ:\r
+        code = 0;\r
+        psrc1 = &src1;\r
+        psrc2 = &src2;\r
+        break;\r
+    case CMP_GE:\r
+        code = 3;\r
+        psrc1 = &src2;\r
+        psrc2 = &src1;\r
+        break;\r
+    case CMP_GT:\r
+        code = 2;\r
+        psrc1 = &src2;\r
+        psrc2 = &src1;\r
+        break;\r
+    case CMP_LE:\r
+        code = 3;\r
+        psrc1 = &src1;\r
+        psrc2 = &src2;\r
+        break;\r
+    case CMP_LT:\r
+        code = 2;\r
+        psrc1 = &src1;\r
+        psrc2 = &src2;\r
+        break;\r
+    case CMP_NE:\r
+        code = 1;\r
+        psrc1 = &src1;\r
+        psrc2 = &src2;\r
+        break;\r
+    default:\r
+        CV_Error(CV_StsBadFlag, "Incorrect compare operation");\r
+    };\r
 \r
-            nppSafeCall( nppiCompare_32f_C1R(src1.ptr<Npp32f>(), static_cast<int>(src1.step),\r
-                src2.ptr<Npp32f>(), static_cast<int>(src2.step),\r
-                dst.ptr<Npp8u>(), static_cast<int>(dst.step), sz, nppCmpOp[cmpop]) );\r
+    dst.create(src1.size(), CV_MAKE_TYPE(CV_8U, src1.channels()));\r
 \r
-            if (stream == 0)\r
-                cudaSafeCall( cudaDeviceSynchronize() );\r
-        }\r
-        else\r
-        {\r
-            device::compare_ne_32f(src1, src2, dst, stream);\r
-        }\r
-    }\r
+    funcs[src1.depth()][code](psrc1->reshape(1), psrc2->reshape(1), dst.reshape(1), StreamAccessor::getStream(stream));\r
 }\r
 \r
 \r

diff --git a/modules/gpu/src/filtering.cpp b/modules/gpu/src/filtering.cpp
index 0f2ba63..747451c 100644 (file)
--- a/modules/gpu/src/filtering.cpp
+++ b/modules/gpu/src/filtering.cpp
@@ -639,17 +639,17 @@ void cv::gpu::morphologyEx(const GpuMat& src, GpuMat& dst, int op, const Mat& ke
     case CV_MOP_GRADIENT:\r
         erode(src, buf2, kernel, buf1, anchor, iterations, stream);\r
         dilate(src, dst, kernel, buf1, anchor, iterations, stream);\r
-        subtract(dst, buf2, dst, stream);\r
+        subtract(dst, buf2, dst, GpuMat(), -1, stream);\r
         break;\r
     case CV_MOP_TOPHAT:\r
         erode(src, dst, kernel, buf1, anchor, iterations, stream);\r
         dilate(dst, buf2, kernel, buf1, anchor, iterations, stream);\r
-        subtract(src, buf2, dst, stream);\r
+        subtract(src, buf2, dst, GpuMat(), -1, stream);\r
         break;\r
     case CV_MOP_BLACKHAT:\r
         dilate(src, dst, kernel, buf1, anchor, iterations, stream);\r
         erode(dst, buf2, kernel, buf1, anchor, iterations, stream);\r
-        subtract(buf2, src, dst, stream);\r
+        subtract(buf2, src, dst, GpuMat(), -1, stream);\r
         break;\r
     default:\r
         CV_Error(CV_StsBadArg, "unknown morphological operation");\r

diff --git a/modules/gpu/test/test_arithm.cpp b/modules/gpu/test/test_arithm.cpp
index 34cadf3..ed9b6cf 100644 (file)
--- a/modules/gpu/test/test_arithm.cpp
+++ b/modules/gpu/test/test_arithm.cpp
@@ -96,7 +96,7 @@ TEST_P(AddArray, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, AddArray, testing::Combine(\r
                         testing::ValuesIn(devices()),\r
-                        testing::Values(CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1)));\r
+                        testing::Values(CV_8UC1, CV_8UC4, CV_16UC1, CV_32SC1, CV_32FC1)));\r
 \r
 struct AddScalar : ArithmTest {};\r
 \r
@@ -130,7 +130,7 @@ TEST_P(AddScalar, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, AddScalar, testing::Combine(\r
                         testing::ValuesIn(devices()),\r
-                        testing::Values(CV_32FC1, CV_32FC2)));\r
+                        testing::Values(CV_8UC1, CV_16UC1, CV_32SC1, CV_32FC1, CV_32FC2)));\r
 \r
 ////////////////////////////////////////////////////////////////////////////////\r
 // subtract\r
@@ -161,7 +161,7 @@ TEST_P(SubtractArray, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, SubtractArray, testing::Combine(\r
                         testing::ValuesIn(devices()),\r
-                        testing::Values(CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1)));\r
+                        testing::Values(CV_8UC1, CV_8UC4, CV_16UC1, CV_32SC1, CV_32FC1)));\r
 \r
 struct SubtractScalar : ArithmTest {};\r
 \r
@@ -195,7 +195,7 @@ TEST_P(SubtractScalar, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, SubtractScalar, testing::Combine(\r
                         testing::ValuesIn(devices()),\r
-                        testing::Values(CV_32FC1, CV_32FC2)));\r
+                        testing::Values(CV_8UC1, CV_16UC1, CV_32SC1, CV_32FC1, CV_32FC2)));\r
 \r
 ////////////////////////////////////////////////////////////////////////////////\r
 // multiply\r
@@ -226,7 +226,7 @@ TEST_P(MultiplyArray, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, MultiplyArray, testing::Combine(\r
                         testing::ValuesIn(devices()),\r
-                        testing::Values(CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1)));\r
+                        testing::Values(CV_8UC1, CV_8UC4, CV_16UC1, CV_32SC1, CV_32FC1)));\r
 \r
 struct MultiplyScalar : ArithmTest {};\r
 \r
@@ -260,7 +260,7 @@ TEST_P(MultiplyScalar, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, MultiplyScalar, testing::Combine(\r
                         testing::ValuesIn(devices()),\r
-                        testing::Values(CV_32FC1)));\r
+                        testing::Values(CV_8UC1, CV_16UC1, CV_32SC1, CV_32FC1)));\r
 \r
 ////////////////////////////////////////////////////////////////////////////////\r
 // divide\r
@@ -291,7 +291,7 @@ TEST_P(DivideArray, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, DivideArray, testing::Combine(\r
                         testing::ValuesIn(devices()),\r
-                        testing::Values(CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1)));\r
+                        testing::Values(CV_8UC1, CV_8UC4, CV_16UC1, CV_32SC1, CV_32FC1)));\r
 \r
 struct DivideScalar : ArithmTest {};\r
 \r
@@ -325,7 +325,7 @@ TEST_P(DivideScalar, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, DivideScalar, testing::Combine(\r
                         testing::ValuesIn(devices()),\r
-                        testing::Values(CV_32FC1)));\r
+                        testing::Values(CV_8UC1, CV_16UC1, CV_32SC1, CV_32FC1)));\r
 \r
 ////////////////////////////////////////////////////////////////////////////////\r
 // transpose\r
@@ -387,7 +387,7 @@ TEST_P(AbsdiffArray, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, AbsdiffArray, testing::Combine(\r
                         testing::ValuesIn(devices()),\r
-                        testing::Values(CV_8UC1, CV_8UC4, CV_32SC1, CV_32FC1)));\r
+                        testing::Values(CV_8UC1, CV_8UC4, CV_16UC1, CV_32SC1, CV_32FC1)));\r
 \r
 struct AbsdiffScalar : ArithmTest {};\r
 \r
@@ -421,7 +421,7 @@ TEST_P(AbsdiffScalar, Accuracy)
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, AbsdiffScalar, testing::Combine(\r
                         testing::ValuesIn(devices()),\r
-                        testing::Values(CV_32FC1)));\r
+                        testing::Values(CV_8UC1, CV_16UC1, CV_32SC1, CV_32FC1)));\r
 \r
 ////////////////////////////////////////////////////////////////////////////////\r
 // compare\r
@@ -813,7 +813,7 @@ TEST_P(Pow, Accuracy)
     /*std::cout  << mat << std::endl << std::endl;\r
     std::cout  << dst << std::endl << std::endl;\r
     std::cout  << dst_gold << std::endl;*/\r
-    EXPECT_MAT_NEAR(dst_gold, dst, 1);\r
+    EXPECT_MAT_NEAR(dst_gold, dst, 2);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(Arithm, Pow, testing::Combine(\r

diff --git a/modules/gpu/test/test_imgproc.cpp b/modules/gpu/test/test_imgproc.cpp
index 493f4a4..30dbb9b 100644 (file)
--- a/modules/gpu/test/test_imgproc.cpp
+++ b/modules/gpu/test/test_imgproc.cpp
@@ -169,8 +169,8 @@ TEST_P(Resize, Accuracy)
         gpuRes2.download(dst2);\r
     );\r
 \r
-    EXPECT_MAT_SIMILAR(dst_gold1, dst1, 0.2);\r
-    EXPECT_MAT_SIMILAR(dst_gold2, dst2, 0.2);\r
+    EXPECT_MAT_SIMILAR(dst_gold1, dst1, 0.21);\r
+    EXPECT_MAT_SIMILAR(dst_gold2, dst2, 0.21);\r
 }\r
 \r
 INSTANTIATE_TEST_CASE_P(ImgProc, Resize, testing::Combine(\r