reduced code convert_to by using templates, merged with copyTo

diff --git a/modules/gpu/src/cuda/matrix_operations.cu b/modules/gpu/src/cuda/matrix_operations.cu
index ad2e0f3..df68b99 100644 (file)
--- a/modules/gpu/src/cuda/matrix_operations.cu
+++ b/modules/gpu/src/cuda/matrix_operations.cu
@@ -53,9 +53,9 @@ __constant__ __align__(16) double scalar_d[4];
 
 namespace mat_operators
 {
-    //////////////////////////////////////////////////////////
-    // CopyTo
-    //////////////////////////////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+////////////////////////////////// CopyTo /////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
 
     template<typename T>
     __global__ void kernel_copy_to_with_mask(T * mat_src, T * mat_dst, const unsigned char * mask, int cols, int rows, int step_mat, int step_mask, int channels)
@@ -71,9 +71,10 @@ namespace mat_operators
             }
     }
 
-    //////////////////////////////////////////////////////////
-    // SetTo
-    //////////////////////////////////////////////////////////
+       
+///////////////////////////////////////////////////////////////////////////
+////////////////////////////////// SetTo //////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
 
     template<typename T>
     __global__ void kernel_set_to_without_mask(T * mat, int cols, int rows, int step, int channels)
@@ -103,156 +104,94 @@ namespace mat_operators
     }
 
 
-    //////////////////////////////////////////////////////////
-    // ConvertTo
-    //////////////////////////////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+//////////////////////////////// ConvertTo ////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
 
-    template <typename T, typename DT, size_t src_elem_size, size_t dst_elem_size>
-    struct Converter
+    template <typename T, typename DT>
+    struct CalcTraits
     {
-        __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
+        __device__ static DT calc(T src, double alpha, double beta)
         {
-            size_t x = threadIdx.x + blockIdx.x * blockDim.x;
-            size_t y = threadIdx.y + blockIdx.y * blockDim.y;
-            if (x < width && y < height)
-            {
-                const T* src = (const T*)(srcmat + src_step * y);
-                DT* dst = (DT*)(dstmat + dst_step * y);
-
-                dst[x] = (DT)__double2int_rn(alpha * src[x] + beta);
-            }
+            return (DT)__double2int_rn(alpha * src + beta);
         }
-        __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
+    };
+    template <typename T>
+    struct CalcTraits<T, float>
+    {
+        __device__ static float calc(T src, double alpha, double beta)
         {
-            return dim3(divUp(width, block.x), divUp(height, block.y));
+            return (float)(alpha * src + beta);
         }
     };
-
-    template <typename T, typename DT>
-    struct Converter<T, DT, 1, 1>
+    template <typename T>
+    struct CalcTraits<T, double>
     {
-        __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
+        __device__ static double calc(T src, double alpha, double beta)
         {
-            size_t x = threadIdx.x + blockIdx.x * blockDim.x;
-            size_t y = threadIdx.y + blockIdx.y * blockDim.y;
-            if (y < height)
-            {
-                const T* src = (const T*)(srcmat + src_step * y);
-                DT* dst = (DT*)(dstmat + dst_step * y);
-                if ((x << 2) + 3 < width)
-                {
-                    uchar4 src4b = ((const uchar4*)src)[x];
-                    uchar4 dst4b;
-
-                    const T* src1b = (const T*) &src4b.x;
-                    DT* dst1b = (DT*) &dst4b.x;
-
-                    dst1b[0] = (DT)__double2int_rn(alpha * src1b[0] + beta);
-                    dst1b[1] = (DT)__double2int_rn(alpha * src1b[1] + beta);
-                    dst1b[2] = (DT)__double2int_rn(alpha * src1b[2] + beta);
-                    dst1b[3] = (DT)__double2int_rn(alpha * src1b[3] + beta);
-
-                    ((uchar4*)dst)[x] = dst4b;
-                }
-                else
-                {
-                    if ((x << 2) + 0 < width)
-                        dst[(x << 2) + 0] = (DT)__double2int_rn(alpha * src[(x << 2) + 0] + beta);
+            return alpha * src + beta;
+        }
+    };
 
-                    if ((x << 2) + 1 < width)
-                        dst[(x << 2) + 1] = (DT)__double2int_rn(alpha * src[(x << 2) + 1] + beta);
+    template <typename T, typename DT, size_t src_elem_size, size_t dst_elem_size>
+    struct ConverterTraits
+    {
+        enum {shift=1};
 
-                    if ((x << 2) + 2 < width)
-                        dst[(x << 2) + 2] = (DT)__double2int_rn(alpha * src[(x << 2) + 2] + beta);
-                }
-            }
-        }
-        __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
-        {
-            return dim3(divUp(width, block.x << 2), divUp(height, block.y));
-        }
-    };/**/
+        typedef T read_type;
+        typedef DT write_type;
+    };
+    template <typename T, typename DT>
+    struct ConverterTraits<T, DT, 1, 1>
+    {
+        enum {shift=4};
 
+        typedef char4 read_type;
+        typedef char4 write_type;
+    };    
     template <typename T, typename DT>
-    struct Converter<T, DT, 1, 2>
+    struct ConverterTraits<T, DT, 2, 1>
     {
-        __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
-        {
-            size_t x = threadIdx.x + blockIdx.x * blockDim.x;
-            size_t y = threadIdx.y + blockIdx.y * blockDim.y;
-            if (y < height)
-            {
-                const T* src = (const T*)(srcmat + src_step * y);
-                DT* dst = (DT*)(dstmat + dst_step * y);
-                if ((x << 1) + 1 < width)
-                {
-                    uchar2 src2b = ((const uchar2*)src)[x];
-                    ushort2 dst2s;
+        enum {shift=4};
 
-                    const T* src1b = (const T*) &src2b;
-                    DT* dst1s = (DT*) &dst2s;
-                    dst1s[0] = (DT)__double2int_rn(alpha * src1b[0] + beta);
-                    dst1s[1] = (DT)__double2int_rn(alpha * src1b[1] + beta);
+        typedef short4 read_type;
+        typedef char4 write_type;
+    };    
+    template <typename T, typename DT>
+    struct ConverterTraits<T, DT, 4, 1>
+    {
+        enum {shift=4};
 
-                    ((ushort2*)(dst))[x] = dst2s;
-                }
-                else
-                {
-                    if ((x << 1) < width)
-                        dst[(x << 1)] = (DT)__double2int_rn(alpha * src[(x << 1)] + beta);
-                }
-            }
-        }
-        __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
-        {
-            return dim3(divUp(width, block.x << 1), divUp(height, block.y));
-        }
-    };/**/
+        typedef int4 read_type;
+        typedef char4 write_type;
+    };    
+    template <typename T, typename DT>
+    struct ConverterTraits<T, DT, 1, 2>
+    {
+        enum {shift=2};
 
+        typedef char2 read_type;
+        typedef short2 write_type;
+    };     
     template <typename T, typename DT>
-    struct Converter<T, DT, 2, 1>
+    struct ConverterTraits<T, DT, 2, 2>
     {
-        __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
-        {
-            size_t x = threadIdx.x + blockIdx.x * blockDim.x;
-            size_t y = threadIdx.y + blockIdx.y * blockDim.y;
-            if (y < height)
-            {
-                const T* src = (const T*)(srcmat + src_step * y);
-                DT* dst = (DT*)(dstmat + dst_step * y);
-                if ((x << 2) + 3 < width)
-                {
-                    ushort4 src4s = ((const ushort4*)src)[x];
-                    uchar4 dst4b;
+        enum {shift=2};
 
-                    const T* src1s = (const T*) &src4s.x;
-                    DT* dst1b = (DT*) &dst4b.x;
-                    dst1b[0] = (DT)__double2int_rn(alpha * src1s[0] + beta);
-                    dst1b[1] = (DT)__double2int_rn(alpha * src1s[1] + beta);
-                    dst1b[2] = (DT)__double2int_rn(alpha * src1s[2] + beta);
-                    dst1b[3] = (DT)__double2int_rn(alpha * src1s[3] + beta);
+        typedef short2 read_type;
+        typedef short2 write_type;
+    };     
+    template <typename T, typename DT>
+    struct ConverterTraits<T, DT, 4, 2>
+    {
+        enum {shift=2};
 
-                    ((uchar4*)(dst))[x] = dst4b;
-                }
-                else
-                {
-                    if ((x << 2) + 0 < width)
-                        dst[(x << 2) + 0] = (DT)__double2int_rn(alpha * src[(x << 2) + 0] + beta);
-                    if ((x << 2) + 1 < width)
-                        dst[(x << 2) + 1] = (DT)__double2int_rn(alpha * src[(x << 2) + 1] + beta);
-                    if ((x << 2) + 2 < width)
-                        dst[(x << 2) + 2] = (DT)__double2int_rn(alpha * src[(x << 2) + 2] + beta);
-                }
-            }
-        }
-        __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
-        {
-            return dim3(divUp(width, block.x << 2), divUp(height, block.y));
-        }
-    };/**/
+        typedef int2 read_type;
+        typedef short2 write_type;
+    };
 
     template <typename T, typename DT>
-    struct Converter<T, DT, 2, 2>
+    struct Converter
     {
         __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
         {
@@ -262,77 +201,37 @@ namespace mat_operators
             {
                 const T* src = (const T*)(srcmat + src_step * y);
                 DT* dst = (DT*)(dstmat + dst_step * y);
-                if ((x << 1) + 1 < width)
+                if ((x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift) + ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift - 1 < width)
                 {
-                    ushort2 src2s = ((const ushort2*)src)[x];
-                    ushort2 dst2s;
+                    typename ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::read_type srcn_el = ((const typename ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::read_type*)src)[x];
+                    typename ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::write_type dstn_el;
 
-                    const T* src1s = (const T*) &src2s.x;
-                    DT* dst1s = (DT*) &dst2s.x;
-                    dst1s[0] = (DT)__double2int_rn(alpha * src1s[0] + beta);
-                    dst1s[1] = (DT)__double2int_rn(alpha * src1s[1] + beta);
+                    const T* src1_el = (const T*) &srcn_el;
+                    DT* dst1_el = (DT*) &dstn_el;
 
-                    ((ushort2*)dst)[x] = dst2s;
+                    for (int i = 0; i < ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift; ++i)
+                        dst1_el[i] = CalcTraits<T, DT>::calc(src1_el[i], alpha, beta);
+
+                    ((typename ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::write_type*)dst)[x] = dstn_el;
                 }
                 else
-                {
-                    if ((x << 1) < width)
-                        dst[(x << 1)] = (DT)__double2int_rn(alpha * src[(x << 1)] + beta);
+                {                    
+                    for (int i = 0; i < ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift - 1; ++i)
+                        if ((x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift) + i < width)
+                            dst[(x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift) + i] = CalcTraits<T, DT>::calc(src[(x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift) + i], alpha, beta);
                 }
             }
         }
         __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
         {
-            return dim3(divUp(width, block.x << 1), divUp(height, block.y));
-        }
-    };/**/
-
-    template <typename T, size_t src_elem_size, size_t dst_elem_size>
-    struct Converter<T, float, src_elem_size, dst_elem_size>
-    {
-        __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
-        {
-            size_t x = threadIdx.x + blockIdx.x * blockDim.x;
-            size_t y = threadIdx.y + blockIdx.y * blockDim.y;
-            if (x < width && y < height)
-            {
-                const T* src = (const T*)(srcmat + src_step * y);
-                float* dst = (float*)(dstmat + dst_step * y);
-
-                dst[x] = (float)(alpha * src[x] + beta);
-            }
-        }
-        __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
-        {
-            return dim3(divUp(width, block.x), divUp(height, block.y));
-        }
-    };
-
-    template <typename T, size_t src_elem_size, size_t dst_elem_size>
-    struct Converter<T, double, src_elem_size, dst_elem_size>
-    {
-        __device__ static void convert(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
-        {
-            size_t x = threadIdx.x + blockIdx.x * blockDim.x;
-            size_t y = threadIdx.y + blockIdx.y * blockDim.y;
-            if (x < width && y < height)
-            {
-                const T* src = (const T*)(srcmat + src_step * y);
-                double* dst = (double*)(dstmat + dst_step * y);
-
-                dst[x] = (double)(alpha * src[x] + beta);
-            }
-        }
-        __host__ static inline dim3 calcGrid(size_t width, size_t height, dim3 block)
-        {
-            return dim3(divUp(width, block.x), divUp(height, block.y));
+            return dim3(divUp(width, block.x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift), divUp(height, block.y));
         }
     };
-
-    template <typename T, typename DT>
+    
+    template <typename T, typename DT> 
     __global__ static void kernel_convert_to(uchar* srcmat, size_t src_step, uchar* dstmat, size_t dst_step, size_t width, size_t height, double alpha, double beta)
     {
-        Converter<T, DT, sizeof(T), sizeof(DT)>::convert(srcmat, src_step, dstmat, dst_step, width, height, alpha, beta);
+        Converter<T, DT>::convert(srcmat, src_step, dstmat, dst_step, width, height, alpha, beta);
     }
 
 } // namespace mat_operators
@@ -344,9 +243,9 @@ namespace cv
                namespace impl
                {
 
-                        //////////////////////////////////////////////////////////////
-                        // CopyTo
-                        //////////////////////////////////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+////////////////////////////////// CopyTo /////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
 
                         typedef void (*CopyToFunc)(const DevMem2D& mat_src, const DevMem2D& mat_dst, const DevMem2D& mask, int channels);
 
@@ -382,9 +281,9 @@ namespace cv
                         }
 
 
-                        //////////////////////////////////////////////////////////////
-                        // SetTo
-                        //////////////////////////////////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+////////////////////////////////// SetTo //////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
 
                         typedef void (*SetToFunc_with_mask)(const DevMem2D& mat, const DevMem2D& mask, int channels);
                         typedef void (*SetToFunc_without_mask)(const DevMem2D& mat, int channels);
@@ -464,59 +363,55 @@ namespace cv
                             func(mat, mask, channels);
                         }
 
-                        //////////////////////////////////////////////////////////////
-                        // ConvertTo
-                        //////////////////////////////////////////////////////////////
-
-                        typedef void (*CvtFunc)(const DevMem2D& src, DevMem2D& dst, size_t width, size_t height, double alpha, double beta);
-
-                        //#if !defined(__CUDA_ARCH__) || (__CUDA_ARCH__ >= 130)
-
-                       template<typename T, typename DT>
-                       void cvt_(const DevMem2D& src, DevMem2D& dst, size_t width, size_t height, double alpha, double beta)
-                       {
-                               dim3 block(32, 8);
-                               dim3 grid = ::mat_operators::Converter<T, DT, sizeof(T), sizeof(DT)>::calcGrid(width, height, block);
-                               ::mat_operators::kernel_convert_to<T, DT><<<grid, block>>>(src.ptr, src.step, dst.ptr, dst.step, width, height, alpha, beta);
-                               cudaSafeCall( cudaThreadSynchronize() );
-                       }
-                       //#endif
-
-                       extern "C" void convert_to(const DevMem2D& src, int sdepth, DevMem2D dst, int ddepth, size_t width, size_t height, double alpha, double beta)
-                       {
-                               static CvtFunc tab[8][8] =
-                               {
-                                       {cvt_<uchar, uchar>, cvt_<uchar, schar>, cvt_<uchar, ushort>, cvt_<uchar, short>,
-                                       cvt_<uchar, int>, cvt_<uchar, float>, cvt_<uchar, double>, 0},
-
-                                       {cvt_<schar, uchar>, cvt_<schar, schar>, cvt_<schar, ushort>, cvt_<schar, short>,
-                                       cvt_<schar, int>, cvt_<schar, float>, cvt_<schar, double>, 0},
-
-                                       {cvt_<ushort, uchar>, cvt_<ushort, schar>, cvt_<ushort, ushort>, cvt_<ushort, short>,
-                                       cvt_<ushort, int>, cvt_<ushort, float>, cvt_<ushort, double>, 0},
-
-                                       {cvt_<short, uchar>, cvt_<short, schar>, cvt_<short, ushort>, cvt_<short, short>,
-                                       cvt_<short, int>, cvt_<short, float>, cvt_<short, double>, 0},
-
-                                       {cvt_<int, uchar>, cvt_<int, schar>, cvt_<int, ushort>,
-                                       cvt_<int, short>, cvt_<int, int>, cvt_<int, float>, cvt_<int, double>, 0},
-
-                                       {cvt_<float, uchar>, cvt_<float, schar>, cvt_<float, ushort>,
-                                       cvt_<float, short>, cvt_<float, int>, cvt_<float, float>, cvt_<float, double>, 0},
-
-                                       {cvt_<double, uchar>, cvt_<double, schar>, cvt_<double, ushort>,
-                                       cvt_<double, short>, cvt_<double, int>, cvt_<double, float>, cvt_<double, double>, 0},
-
-                                       {0,0,0,0,0,0,0,0}
-                               };
-
-                               CvtFunc func = tab[sdepth][ddepth];
-                               if (func == 0)
-                                       error("Operation \'ConvertTo\' doesn't supported on your GPU model", __FILE__, __LINE__);
-                               func(src, dst, width, height, alpha, beta);
-                       }
-               }
-
-
-       }
-}
+                                               
+///////////////////////////////////////////////////////////////////////////
+//////////////////////////////// ConvertTo ////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+
+                                   typedef void (*CvtFunc)(const DevMem2D& src, DevMem2D& dst, size_t width, size_t height, double alpha, double beta);
+
+                                   template<typename T, typename DT> 
+                                   void cvt_(const DevMem2D& src, DevMem2D& dst, size_t width, size_t height, double alpha, double beta)
+                                   {
+                                           dim3 block(32, 8);
+                                           dim3 grid = ::mat_operators::Converter<T, DT>::calcGrid(width, height, block);
+                                           ::mat_operators::kernel_convert_to<T, DT><<<grid, block>>>(src.ptr, src.step, dst.ptr, dst.step, width, height, alpha, beta);
+                                           cudaSafeCall( cudaThreadSynchronize() );
+                                   }
+
+                                   extern "C" void convert_to(const DevMem2D& src, int sdepth, DevMem2D dst, int ddepth, size_t width, size_t height, double alpha, double beta)
+                                   {
+                                           static CvtFunc tab[8][8] =
+                                           {
+                                                   {cvt_<uchar, uchar>, cvt_<uchar, schar>, cvt_<uchar, ushort>, cvt_<uchar, short>,
+                                                   cvt_<uchar, int>, cvt_<uchar, float>, cvt_<uchar, double>, 0},
+
+                                                   {cvt_<schar, uchar>, cvt_<schar, schar>, cvt_<schar, ushort>, cvt_<schar, short>,
+                                                   cvt_<schar, int>, cvt_<schar, float>, cvt_<schar, double>, 0},
+
+                                                   {cvt_<ushort, uchar>, cvt_<ushort, schar>, cvt_<ushort, ushort>, cvt_<ushort, short>,
+                                                   cvt_<ushort, int>, cvt_<ushort, float>, cvt_<ushort, double>, 0},
+
+                                                   {cvt_<short, uchar>, cvt_<short, schar>, cvt_<short, ushort>, cvt_<short, short>,
+                                                   cvt_<short, int>, cvt_<short, float>, cvt_<short, double>, 0},
+
+                                                   {cvt_<int, uchar>, cvt_<int, schar>, cvt_<int, ushort>,
+                                                   cvt_<int, short>, cvt_<int, int>, cvt_<int, float>, cvt_<int, double>, 0},
+
+                                                   {cvt_<float, uchar>, cvt_<float, schar>, cvt_<float, ushort>,
+                                                   cvt_<float, short>, cvt_<float, int>, cvt_<float, float>, cvt_<float, double>, 0},
+
+                                                   {cvt_<double, uchar>, cvt_<double, schar>, cvt_<double, ushort>,
+                                                   cvt_<double, short>, cvt_<double, int>, cvt_<double, float>, cvt_<double, double>, 0},
+
+                                                   {0,0,0,0,0,0,0,0}
+                                           };
+
+                                           CvtFunc func = tab[sdepth][ddepth];
+                                           if (func == 0)
+                                cv::gpu::error("Operation \'ConvertTo\' doesn't supported on your GPU model", __FILE__, __LINE__);
+                                           func(src, dst, width, height, alpha, beta);
+                                   }
+               } // namespace impl             
+       } // namespace gpu
+} // namespace cv

diff --git a/modules/gpu/src/matrix_operations.cpp b/modules/gpu/src/matrix_operations.cpp
index 9469f59..f64733d 100644 (file)
--- a/modules/gpu/src/matrix_operations.cpp
+++ b/modules/gpu/src/matrix_operations.cpp
@@ -114,8 +114,6 @@ void cv::gpu::GpuMat::copyTo( GpuMat& mat, const GpuMat& mask ) const
 \r
 void cv::gpu::GpuMat::convertTo( GpuMat& dst, int rtype, double alpha, double beta ) const\r
 {\r
-    //CV_Assert(!"Not implemented");\r
-\r
     bool noScale = fabs(alpha-1) < std::numeric_limits<double>::epsilon() && fabs(beta) < std::numeric_limits<double>::epsilon();\r
 \r
     if( rtype < 0 )\r
@@ -124,11 +122,11 @@ void cv::gpu::GpuMat::convertTo( GpuMat& dst, int rtype, double alpha, double be
         rtype = CV_MAKETYPE(CV_MAT_DEPTH(rtype), channels());\r
 \r
     int sdepth = depth(), ddepth = CV_MAT_DEPTH(rtype);\r
-    /*if( sdepth == ddepth && noScale )\r
+    if( sdepth == ddepth && noScale )\r
     {\r
         copyTo(dst);\r
         return;\r
-    }*/\r
+    }\r
 \r
     GpuMat temp;\r
     const GpuMat* psrc = this;\r

diff --git a/tests/gpu/src/convert_to.cpp b/tests/gpu/src/convert_to.cpp
index c53fbc4..b2e231a 100644 (file)
--- a/tests/gpu/src/convert_to.cpp
+++ b/tests/gpu/src/convert_to.cpp
@@ -26,7 +26,7 @@ void CV_GpuMatOpConvertTo::run( int /* start_from */)
 {
     const Size img_size(67, 35);
 
-    const int types[] = {CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F/**/};
+    const int types[] = {CV_8U, CV_8S, CV_16U, CV_16S, CV_32S, CV_32F, CV_64F};
     const int types_num = sizeof(types) / sizeof(int);
 
     const char* types_str[] = {"CV_8U", "CV_8S", "CV_16U", "CV_16S", "CV_32S", "CV_32F", "CV_64F"};
@@ -39,9 +39,6 @@ void CV_GpuMatOpConvertTo::run( int /* start_from */)
         {
             for (int c = 1; c < 2 && passed; ++c)
             {
-                //if (i == j)
-                  //  continue;
-
                 const int src_type = CV_MAKETYPE(types[i], c);
                 const int dst_type = types[j];
                 const double alpha = (double)rand() / RAND_MAX * 10.0;
@@ -53,30 +50,34 @@ void CV_GpuMatOpConvertTo::run( int /* start_from */)
                 Mat cpumatdst;
                 GpuMat gpumatdst;
                 
-                //double cput = (double)getTickCount();
-                cpumatsrc.convertTo(cpumatdst, dst_type, alpha, beta);
-                //cput = ((double)getTickCount() - cput) / getTickFrequency();
+                //TickMeter tm;
+                //tm.start();
+                //for(int i = 0; i < 50; ++i)
+                    cpumatsrc.convertTo(cpumatdst, dst_type, alpha, beta);
+                //tm.stop();
+                //cout << "SRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << endl << "\tCPU FPS = " << 50.0/tm.getTimeSec() << endl;
 
-                //double gput = (double)getTickCount();
-                gpumatsrc.convertTo(gpumatdst, dst_type, alpha, beta);
-                //gput = ((double)getTickCount() - gput) / getTickFrequency();
+                //tm.reset();
 
-                /*cout << "convertTo time: " << endl;
-                cout << "CPU time: " << cput << endl;
-                cout << "GPU time: " << gput << endl;/**/               
+                try
+                {
+                    //tm.start();
+                    //for(int i = 0; i < 50; ++i)
+                        gpumatsrc.convertTo(gpumatdst, dst_type, alpha, beta);
+                    //tm.stop();
+                    //cout << "\tGPU FPS = " << 50.0/tm.getTimeSec() << endl;
+                }
+                catch(cv::Exception& e)
+                {
+                    cout << "ERROR: " << e.err << endl;
+                    passed = false;
+                    break;
+                }
 
-                double r = norm(cpumatdst, gpumatdst, NORM_L1);
+                double r = norm(cpumatdst, gpumatdst, NORM_INF);
                 if (r > 1)
-                {
-                    /*namedWindow("CPU"); \r
-                    imshow("CPU", cpumatdst); 
-                    namedWindow("GPU"); \r
-                    imshow("GPU", gpumatdst); \r
-                    waitKey();/**/
-                    
-                    cout << "Failed:" << endl;
-                    cout << "\tr = " << r << endl;
-                    cout << "\tSRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << endl;/**/
+                {                    
+                    cout << "FAILED: " << "SRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << " NORM = " << r << endl;
 
                     passed = false;
                 }