diff --git a/modules/core/src/cuda/matrix_operations.cu b/modules/core/src/cuda/matrix_operations.cu index 9e830e563..60aa07340 100644 --- a/modules/core/src/cuda/matrix_operations.cu +++ b/modules/core/src/cuda/matrix_operations.cu @@ -44,6 +44,7 @@ #include "opencv2/gpu/device/saturate_cast.hpp" #include "opencv2/gpu/device/transform.hpp" #include "opencv2/gpu/device/functional.hpp" +#include "opencv2/gpu/device/type_traits.hpp" namespace cv { namespace gpu { namespace device { @@ -54,6 +55,7 @@ namespace cv { namespace gpu { namespace device void writeScalar(const int*); void writeScalar(const float*); void writeScalar(const double*); + void copyToWithMask_gpu(PtrStepSzb src, PtrStepSzb dst, size_t elemSize1, int cn, PtrStepSzb mask, bool colorMask, cudaStream_t stream); void convert_gpu(PtrStepSzb, int, PtrStepSzb, int, double, double, cudaStream_t); }}} @@ -226,16 +228,16 @@ namespace cv { namespace gpu { namespace device //////////////////////////////// ConvertTo //////////////////////////////// /////////////////////////////////////////////////////////////////////////// - template struct Convertor : unary_function + template struct Convertor : unary_function { - Convertor(double alpha_, double beta_) : alpha(alpha_), beta(beta_) {} + Convertor(S alpha_, S beta_) : alpha(alpha_), beta(beta_) {} - __device__ __forceinline__ D operator()(const T& src) const + __device__ __forceinline__ D operator()(typename TypeTraits::ParameterType src) const { return saturate_cast(alpha * src + beta); } - double alpha, beta; + S alpha, beta; }; namespace detail @@ -282,16 +284,16 @@ namespace cv { namespace gpu { namespace device }; } - template struct TransformFunctorTraits< Convertor > : detail::ConvertTraits< Convertor > + template struct TransformFunctorTraits< Convertor > : detail::ConvertTraits< Convertor > { }; - template + template void cvt_(PtrStepSzb src, PtrStepSzb dst, double alpha, double beta, cudaStream_t stream) { cudaSafeCall( cudaSetDoubleForDevice(&alpha) ); cudaSafeCall( cudaSetDoubleForDevice(&beta) ); - Convertor op(alpha, beta); + Convertor op(static_cast(alpha), static_cast(beta)); cv::gpu::device::transform((PtrStepSz)src, (PtrStepSz)dst, op, WithOutMask(), stream); } @@ -304,36 +306,74 @@ namespace cv { namespace gpu { namespace device { typedef void (*caller_t)(PtrStepSzb src, PtrStepSzb dst, double alpha, double beta, cudaStream_t stream); - static const caller_t tab[8][8] = + static const caller_t tab[7][7] = { - {cvt_, cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, 0}, - - {cvt_, cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, 0}, - - {cvt_, cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, 0}, - - {cvt_, cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, 0}, - - {cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, cvt_, 0}, - - {cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, cvt_, 0}, - - {cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, cvt_, 0}, - - {0,0,0,0,0,0,0,0} + { + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_ + }, + { + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_ + }, + { + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_ + }, + { + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_ + }, + { + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_ + }, + { + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_ + }, + { + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_, + cvt_ + } }; caller_t func = tab[sdepth][ddepth]; - if (!func) - cv::gpu::error("Unsupported convert operation", __FILE__, __LINE__, "convert_gpu"); - func(src, dst, alpha, beta, stream); }