diff --git a/modules/gpu/src/cuda/matrix_operations.cu b/modules/gpu/src/cuda/matrix_operations.cu index 1a7fe69d81..84c029d944 100644 --- a/modules/gpu/src/cuda/matrix_operations.cu +++ b/modules/gpu/src/cuda/matrix_operations.cu @@ -44,6 +44,7 @@ #include #include "cuda_shared.hpp" #include "cuda_runtime.h" +#include "saturate_cast.hpp" using namespace cv::gpu; using namespace cv::gpu::impl; @@ -108,31 +109,6 @@ namespace mat_operators //////////////////////////////// ConvertTo //////////////////////////////// /////////////////////////////////////////////////////////////////////////// - template - struct ScaleTraits - { - __device__ static DT scale(T src, double alpha, double beta) - { - return (DT)__double2int_rn(alpha * src + beta); - } - }; - template - struct ScaleTraits - { - __device__ static float scale(T src, double alpha, double beta) - { - return (float)(alpha * src + beta); - } - }; - template - struct ScaleTraits - { - __device__ static double scale(T src, double alpha, double beta) - { - return alpha * src + beta; - } - }; - template struct ReadWriteTraits { @@ -213,7 +189,7 @@ namespace mat_operators DT* dst1_el = (DT*) &dstn_el; for (int i = 0; i < shift; ++i) - dst1_el[i] = ScaleTraits::scale(src1_el[i], alpha, beta); + dst1_el[i] = saturate_cast
(alpha * src1_el[i] + beta); ((write_type*)dst)[x] = dstn_el; } @@ -221,7 +197,7 @@ namespace mat_operators { for (int i = 0; i < shift - 1; ++i) if ((x * shift) + i < width) - dst[(x * shift) + i] = ScaleTraits::scale(src[(x * shift) + i], alpha, beta); + dst[(x * shift) + i] = saturate_cast
(alpha * src[(x * shift) + i] + beta); } } } diff --git a/modules/gpu/src/cuda/saturate_cast.hpp b/modules/gpu/src/cuda/saturate_cast.hpp index 027ea29386..c8fbc7ebb3 100644 --- a/modules/gpu/src/cuda/saturate_cast.hpp +++ b/modules/gpu/src/cuda/saturate_cast.hpp @@ -43,11 +43,127 @@ #ifndef __OPENCV_GPU_SATURATE_CAST_HPP__ #define __OPENCV_GPU_SATURATE_CAST_HPP__ +#include "cuda_shared.hpp" -template -__device__ void saturate_cast(F) +namespace cv { - + namespace gpu + { + template __device__ _Tp saturate_cast(uchar v) { return _Tp(v); } + template __device__ _Tp saturate_cast(schar v) { return _Tp(v); } + template __device__ _Tp saturate_cast(ushort v) { return _Tp(v); } + template __device__ _Tp saturate_cast(short v) { return _Tp(v); } + template __device__ _Tp saturate_cast(uint v) { return _Tp(v); } + template __device__ _Tp saturate_cast(int v) { return _Tp(v); } + template __device__ _Tp saturate_cast(float v) { return _Tp(v); } + template __device__ _Tp saturate_cast(double v) { return _Tp(v); } + + template<> __device__ uchar saturate_cast(schar v) + { return (uchar)max((int)v, 0); } + template<> __device__ uchar saturate_cast(ushort v) + { return (uchar)min((uint)v, (uint)UCHAR_MAX); } + template<> __device__ uchar saturate_cast(int v) + { return (uchar)((uint)v <= UCHAR_MAX ? v : v > 0 ? UCHAR_MAX : 0); } + template<> __device__ uchar saturate_cast(uint v) + { return (uchar)min(v, (uint)UCHAR_MAX); } + template<> __device__ uchar saturate_cast(short v) + { return saturate_cast((uint)v); } + + template<> __device__ uchar saturate_cast(float v) + { int iv = __float2int_rn(v); return saturate_cast(iv); } + template<> __device__ uchar saturate_cast(double v) + { + #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130 + int iv = __double2int_rn(v); return saturate_cast(iv); + #else + return saturate_cast((float)v); + #endif + } + + template<> __device__ schar saturate_cast(uchar v) + { return (schar)min((int)v, SCHAR_MAX); } + template<> __device__ schar saturate_cast(ushort v) + { return (schar)min((uint)v, (uint)SCHAR_MAX); } + template<> __device__ schar saturate_cast(int v) + { + return (schar)((uint)(v-SCHAR_MIN) <= (uint)UCHAR_MAX ? + v : v > 0 ? SCHAR_MAX : SCHAR_MIN); + } + template<> __device__ schar saturate_cast(short v) + { return saturate_cast((int)v); } + template<> __device__ schar saturate_cast(uint v) + { return (schar)min(v, (uint)SCHAR_MAX); } + + template<> __device__ schar saturate_cast(float v) + { int iv = __float2int_rn(v); return saturate_cast(iv); } + template<> __device__ schar saturate_cast(double v) + { + #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130 + int iv = __double2int_rn(v); return saturate_cast(iv); + #else + return saturate_cast((float)v); + #endif + } + + template<> __device__ ushort saturate_cast(schar v) + { return (ushort)max((int)v, 0); } + template<> __device__ ushort saturate_cast(short v) + { return (ushort)max((int)v, 0); } + template<> __device__ ushort saturate_cast(int v) + { return (ushort)((uint)v <= (uint)USHRT_MAX ? v : v > 0 ? USHRT_MAX : 0); } + template<> __device__ ushort saturate_cast(uint v) + { return (ushort)min(v, (uint)USHRT_MAX); } + template<> __device__ ushort saturate_cast(float v) + { int iv = __float2int_rn(v); return saturate_cast(iv); } + template<> __device__ ushort saturate_cast(double v) + { + #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130 + int iv = __double2int_rn(v); return saturate_cast(iv); + #else + return saturate_cast((float)v); + #endif + } + + template<> __device__ short saturate_cast(ushort v) + { return (short)min((int)v, SHRT_MAX); } + template<> __device__ short saturate_cast(int v) + { + return (short)((uint)(v - SHRT_MIN) <= (uint)USHRT_MAX ? + v : v > 0 ? SHRT_MAX : SHRT_MIN); + } + template<> __device__ short saturate_cast(uint v) + { return (short)min(v, (uint)SHRT_MAX); } + template<> __device__ short saturate_cast(float v) + { int iv = __float2int_rn(v); return saturate_cast(iv); } + template<> __device__ short saturate_cast(double v) + { + #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130 + int iv = __double2int_rn(v); return saturate_cast(iv); + #else + return saturate_cast((float)v); + #endif + } + + template<> __device__ int saturate_cast(float v) { return __float2int_rn(v); } + template<> __device__ int saturate_cast(double v) + { + #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130 + return __double2int_rn(v); + #else + return saturate_cast((float)v); + #endif + } + + template<> __device__ uint saturate_cast(float v){ return __float2uint_rn(v); } + template<> __device__ uint saturate_cast(double v) + { + #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130 + return __double2uint_rn(v); + #else + return saturate_cast((float)v); + #endif + } + } } #endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */ \ No newline at end of file diff --git a/tests/gpu/src/convert_to.cpp b/tests/gpu/src/convert_to.cpp index e4b205ec50..4bebba6f20 100644 --- a/tests/gpu/src/convert_to.cpp +++ b/tests/gpu/src/convert_to.cpp @@ -32,51 +32,47 @@ void CV_GpuMatOpConvertTo::run( int /* start_from */) const char* types_str[] = {"CV_8U", "CV_8S", "CV_16U", "CV_16S", "CV_32S", "CV_32F", "CV_64F"}; bool passed = true; - - for (int i = 0; i < types_num && passed; ++i) + try { - for (int j = 0; j < types_num && passed; ++j) + for (int i = 0; i < types_num && passed; ++i) { - for (int c = 1; c < 2 && passed; ++c) + for (int j = 0; j < types_num && passed; ++j) { - const int src_type = CV_MAKETYPE(types[i], c); - const int dst_type = types[j]; - const double alpha = (double)rand() / RAND_MAX * 10.0; - const double beta = (double)rand() / RAND_MAX * 10.0; - - cv::RNG rng(*ts->get_rng()); - - Mat cpumatsrc(img_size, src_type); - - rng.fill(cpumatsrc, RNG::UNIFORM, Scalar::all(0), Scalar::all(10)); - - GpuMat gpumatsrc(cpumatsrc); - Mat cpumatdst; - GpuMat gpumatdst; - - cpumatsrc.convertTo(cpumatdst, dst_type, alpha, beta); - - try + for (int c = 1; c < 2 && passed; ++c) { + const int src_type = CV_MAKETYPE(types[i], c); + const int dst_type = types[j]; + const double alpha = (double)rand() / RAND_MAX * 2.0; + const double beta = (double)rand() / RAND_MAX * 150.0 - 75; + + cv::RNG rng(*ts->get_rng()); + + Mat cpumatsrc(img_size, src_type); + + rng.fill(cpumatsrc, RNG::UNIFORM, Scalar::all(0), Scalar::all(300)); + + GpuMat gpumatsrc(cpumatsrc); + Mat cpumatdst; + GpuMat gpumatdst; + + cpumatsrc.convertTo(cpumatdst, dst_type, alpha, beta); gpumatsrc.convertTo(gpumatdst, dst_type, alpha, beta); - } - catch(cv::Exception& e) - { - cout << "ERROR: " << e.err << endl; - passed = false; - break; - } + + double r = norm(cpumatdst, gpumatdst, NORM_INF); + if (r > 1) + { + cout << "FAILED: " << "SRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << " NORM = " << r << endl; - double r = norm(cpumatdst, gpumatdst, NORM_INF); - if (r > 1) - { - cout << "FAILED: " << "SRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << " NORM = " << r << endl; - - passed = false; + passed = false; + } } } } } + catch(cv::Exception& e) + { + cout << "ERROR: " << e.err << endl; + } ts->set_failed_test_info(passed ? CvTS::OK : CvTS::FAIL_GENERIC); }