\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
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
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
\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
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
\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
\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
\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