implemented saturate_cast on gpu; updated GpuMat::convertTo implementation with saturate_cast; minor fix of convert_to test

modules/gpu/src/cuda/matrix_operations.cu

@@ -44,6 +44,7 @@
 #include <stdio.h>
 #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 <typename T, typename DT>
-    struct ScaleTraits
-    {
-        __device__ static DT scale(T src, double alpha, double beta)
-        {
-            return (DT)__double2int_rn(alpha * src + beta);
-        }
-    };
-    template <typename T>
-    struct ScaleTraits<T, float>
-    {
-        __device__ static float scale(T src, double alpha, double beta)
-        {
-            return (float)(alpha * src + beta);
-        }
-    };
-    template <typename T>
-    struct ScaleTraits<T, double>
-    {
-        __device__ static double scale(T src, double alpha, double beta)
-        {
-            return alpha * src + beta;
-        }
-    };
-
     template <typename T, typename DT, size_t src_elem_size, size_t dst_elem_size>
     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<T, DT>::scale(src1_el[i], alpha, beta);
+                    dst1_el[i] =  saturate_cast<DT>(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<T, DT>::scale(src[(x * shift) + i], alpha, beta);
+                        dst[(x * shift) + i] = saturate_cast<DT>(alpha * src[(x * shift) + i] + beta);
             }
         }
     }

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<class F,T>
+namespace cv
-__device__ void saturate_cast<T>(F)
 {
-
+    namespace gpu
+    {
+        template<typename _Tp> __device__ _Tp saturate_cast(uchar v) { return _Tp(v); }
+        template<typename _Tp> __device__ _Tp saturate_cast(schar v) { return _Tp(v); }
+        template<typename _Tp> __device__ _Tp saturate_cast(ushort v) { return _Tp(v); }
+        template<typename _Tp> __device__ _Tp saturate_cast(short v) { return _Tp(v); }
+        template<typename _Tp> __device__ _Tp saturate_cast(uint v) { return _Tp(v); }
+        template<typename _Tp> __device__ _Tp saturate_cast(int v) { return _Tp(v); }
+        template<typename _Tp> __device__ _Tp saturate_cast(float v) { return _Tp(v); }
+        template<typename _Tp> __device__ _Tp saturate_cast(double v) { return _Tp(v); }
+
+        template<> __device__ uchar saturate_cast<uchar>(schar v)
+        { return (uchar)max((int)v, 0); }
+        template<> __device__ uchar saturate_cast<uchar>(ushort v)
+        { return (uchar)min((uint)v, (uint)UCHAR_MAX); }
+        template<> __device__ uchar saturate_cast<uchar>(int v)
+        { return (uchar)((uint)v <= UCHAR_MAX ? v : v > 0 ? UCHAR_MAX : 0); }
+        template<> __device__ uchar saturate_cast<uchar>(uint v)
+        { return (uchar)min(v, (uint)UCHAR_MAX); }
+        template<> __device__ uchar saturate_cast<uchar>(short v)
+        { return saturate_cast<uchar>((uint)v); }
+
+        template<> __device__ uchar saturate_cast<uchar>(float v)
+        { int iv = __float2int_rn(v); return saturate_cast<uchar>(iv); }
+        template<> __device__ uchar saturate_cast<uchar>(double v)
+        {
+        #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
+            int iv = __double2int_rn(v); return saturate_cast<uchar>(iv);
+        #else
+            return saturate_cast<uchar>((float)v);
+        #endif
+        }
+
+        template<> __device__ schar saturate_cast<schar>(uchar v)
+        { return (schar)min((int)v, SCHAR_MAX); }
+        template<> __device__ schar saturate_cast<schar>(ushort v)
+        { return (schar)min((uint)v, (uint)SCHAR_MAX); }
+        template<> __device__ schar saturate_cast<schar>(int v)
+        {
+            return (schar)((uint)(v-SCHAR_MIN) <= (uint)UCHAR_MAX ?
+                        v : v > 0 ? SCHAR_MAX : SCHAR_MIN);
+        }
+        template<> __device__ schar saturate_cast<schar>(short v)
+        { return saturate_cast<schar>((int)v); }
+        template<> __device__ schar saturate_cast<schar>(uint v)
+        { return (schar)min(v, (uint)SCHAR_MAX); }
+
+        template<> __device__ schar saturate_cast<schar>(float v)
+        { int iv = __float2int_rn(v); return saturate_cast<schar>(iv); }
+        template<> __device__ schar saturate_cast<schar>(double v)
+        {             
+        #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
+            int iv = __double2int_rn(v); return saturate_cast<schar>(iv);
+        #else
+            return saturate_cast<schar>((float)v);
+        #endif
+        }
+
+        template<> __device__ ushort saturate_cast<ushort>(schar v)
+        { return (ushort)max((int)v, 0); }
+        template<> __device__ ushort saturate_cast<ushort>(short v)
+        { return (ushort)max((int)v, 0); }
+        template<> __device__ ushort saturate_cast<ushort>(int v)
+        { return (ushort)((uint)v <= (uint)USHRT_MAX ? v : v > 0 ? USHRT_MAX : 0); }
+        template<> __device__ ushort saturate_cast<ushort>(uint v)
+        { return (ushort)min(v, (uint)USHRT_MAX); }
+        template<> __device__ ushort saturate_cast<ushort>(float v)
+        { int iv = __float2int_rn(v); return saturate_cast<ushort>(iv); }
+        template<> __device__ ushort saturate_cast<ushort>(double v)
+        {             
+        #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
+            int iv = __double2int_rn(v); return saturate_cast<ushort>(iv);
+        #else
+            return saturate_cast<ushort>((float)v);
+        #endif
+        }
+
+        template<> __device__ short saturate_cast<short>(ushort v)
+        { return (short)min((int)v, SHRT_MAX); }
+        template<> __device__ short saturate_cast<short>(int v)
+        {
+            return (short)((uint)(v - SHRT_MIN) <= (uint)USHRT_MAX ?
+                    v : v > 0 ? SHRT_MAX : SHRT_MIN);
+        }
+        template<> __device__ short saturate_cast<short>(uint v)
+        { return (short)min(v, (uint)SHRT_MAX); }
+        template<> __device__ short saturate_cast<short>(float v)
+        { int iv = __float2int_rn(v); return saturate_cast<short>(iv); }
+        template<> __device__ short saturate_cast<short>(double v)
+        {            
+        #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
+            int iv = __double2int_rn(v); return saturate_cast<short>(iv);
+        #else
+            return saturate_cast<short>((float)v);
+        #endif
+        }
+
+        template<> __device__ int saturate_cast<int>(float v) { return __float2int_rn(v); }
+        template<> __device__ int saturate_cast<int>(double v) 
+        {
+        #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130 
+            return __double2int_rn(v);
+        #else
+            return saturate_cast<int>((float)v);
+        #endif
+        }
+
+        template<> __device__ uint saturate_cast<uint>(float v){ return __float2uint_rn(v); }
+        template<> __device__ uint saturate_cast<uint>(double v) 
+        {            
+        #if defined __CUDA_ARCH__ && __CUDA_ARCH__ >= 130
+            return __double2uint_rn(v);
+        #else
+            return saturate_cast<uint>((float)v);
+        #endif
+        }
+    }
 }
 
 #endif /* __OPENCV_GPU_SATURATE_CAST_HPP__ */

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;
-
+    try
-    for (int i = 0; i < types_num && passed; ++i)
     {
-        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);
+                for (int c = 1; c < 2 && passed; ++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
                 {
+                    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)
+                    double r = norm(cpumatdst, gpumatdst, NORM_INF);
-                {
+                    if (r > 1)
-                    cout << "ERROR: " << e.err << endl;
+                    {
-                    passed = false;
+                        cout << "FAILED: " << "SRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << " NORM = " << r << endl;
-                    break;
-                }
 
-                double r = norm(cpumatdst, gpumatdst, NORM_INF);
+                        passed = false;
-                if (r > 1)
+                    }
-                {
-                    cout << "FAILED: " << "SRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << " NORM = " << r << endl;
-
-                    passed = false;
                 }
             }
         }
     }
+    catch(cv::Exception& e)
+    {
+        cout << "ERROR: " << e.err << endl;
+    }
     ts->set_failed_test_info(passed ? CvTS::OK : CvTS::FAIL_GENERIC);
 }