reduced code convert_to by using templates, merged with copyTo

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,33 +104,94 @@ namespace mat_operators
     }
 
 
-    //////////////////////////////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
-    // ConvertTo
+//////////////////////////////// ConvertTo ////////////////////////////////
-    //////////////////////////////////////////////////////////
+///////////////////////////////////////////////////////////////////////////
+
+    template <typename T, typename DT>
+    struct CalcTraits
+    {
+        __device__ static DT calc(T src, double alpha, double beta)
+        {
+            return (DT)__double2int_rn(alpha * src + beta);
+        }
+    };
+    template <typename T>
+    struct CalcTraits<T, float>
+    {
+        __device__ static float calc(T src, double alpha, double beta)
+        {
+            return (float)(alpha * src + beta);
+        }
+    };
+    template <typename T>
+    struct CalcTraits<T, double>
+    {
+        __device__ static double calc(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 ConverterTraits
+    {
+        enum {shift=1};
+
+        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 ConverterTraits<T, DT, 2, 1>
+    {
+        enum {shift=4};
+
+        typedef short4 read_type;
+        typedef char4 write_type;
+    };    
+    template <typename T, typename DT>
+    struct ConverterTraits<T, DT, 4, 1>
+    {
+        enum {shift=4};
+
+        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 ConverterTraits<T, DT, 2, 2>
+    {
+        enum {shift=2};
+
+        typedef short2 read_type;
+        typedef short2 write_type;
+    };     
+    template <typename T, typename DT>
+    struct ConverterTraits<T, DT, 4, 2>
+    {
+        enum {shift=2};
+
+        typedef int2 read_type;
+        typedef short2 write_type;
+    };
+
+    template <typename T, typename DT>
     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)
-        {
-            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);
-            }
-        }
-        __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, typename DT>
-    struct Converter<T, DT, 1, 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)
         {
@@ -139,200 +201,37 @@ namespace mat_operators
             {
                 const T* src = (const T*)(srcmat + src_step * y);
                 DT* dst = (DT*)(dstmat + dst_step * y);
-                if ((x << 2) + 3 < width)
+                if ((x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift) + ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift - 1 < width)
                 {
-                    uchar4 src4b = ((const uchar4*)src)[x];
+                    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];
-                    uchar4 dst4b;
+                    typename ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::write_type dstn_el;
 
-                    const T* src1b = (const T*) &src4b.x;
+                    const T* src1_el = (const T*) &srcn_el;
-                    DT* dst1b = (DT*) &dst4b.x;
+                    DT* dst1_el = (DT*) &dstn_el;
 
-                    dst1b[0] = (DT)__double2int_rn(alpha * src1b[0] + beta);
+                    for (int i = 0; i < ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift; ++i)
-                    dst1b[1] = (DT)__double2int_rn(alpha * src1b[1] + beta);
+                        dst1_el[i] = CalcTraits<T, DT>::calc(src1_el[i], alpha, beta);
-                    dst1b[2] = (DT)__double2int_rn(alpha * src1b[2] + beta);
-                    dst1b[3] = (DT)__double2int_rn(alpha * src1b[3] + beta);
 
-                    ((uchar4*)dst)[x] = dst4b;
+                    ((typename ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::write_type*)dst)[x] = dstn_el;
                 }
                 else
                 {                    
-                    if ((x << 2) + 0 < width)
+                    for (int i = 0; i < ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift - 1; ++i)
-                        dst[(x << 2) + 0] = (DT)__double2int_rn(alpha * src[(x << 2) + 0] + beta);
+                        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);
-                    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));
+            return dim3(divUp(width, block.x * ConverterTraits<T, DT, sizeof(T), sizeof(DT)>::shift), divUp(height, block.y));
-        }
-    };/**/
-
-    template <typename T, typename DT>
-    struct Converter<T, DT, 1, 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 << 1) + 1 < width)
-                {
-                    uchar2 src2b = ((const uchar2*)src)[x];
-                    ushort2 dst2s;
-
-                    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);
-
-                    ((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));
-        }
-    };/**/
-
-    template <typename T, typename DT>
-    struct Converter<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 << 2) + 3 < width)
-                {
-                    ushort4 src4s = ((const ushort4*)src)[x];
-                    uchar4 dst4b;
-
-                    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);
-
-                    ((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));
-        }
-    };/**/
-
-    template <typename T, typename DT>
-    struct Converter<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 << 1) + 1 < width)
-                {
-                    ushort2 src2s = ((const ushort2*)src)[x];
-                    ushort2 dst2s;
-
-                    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);
-
-                    ((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));
-        }
-    };/**/
-
-    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));
         }
     };
     
     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,23 +363,21 @@ namespace cv
                             func(mat, mask, channels);
                         }
 
-                        //////////////////////////////////////////////////////////////
+						
-                        // ConvertTo
+///////////////////////////////////////////////////////////////////////////
-                        //////////////////////////////////////////////////////////////
+//////////////////////////////// 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);
+				            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() );
 			            }
-			//#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)
 			            {
@@ -512,11 +409,9 @@ namespace cv
 
 				            CvtFunc func = tab[sdepth][ddepth];
 				            if (func == 0)
-					error("Operation \'ConvertTo\' doesn't supported on your GPU model", __FILE__, __LINE__);
+                                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
-	}
-}

modules/gpu/src/matrix_operations.cpp

@@ -114,8 +114,6 @@ void cv::gpu::GpuMat::copyTo( GpuMat& mat, const GpuMat& mask ) const
 
 void cv::gpu::GpuMat::convertTo( GpuMat& dst, int rtype, double alpha, double beta ) const
 {
-    //CV_Assert(!"Not implemented");
-
     bool noScale = fabs(alpha-1) < std::numeric_limits<double>::epsilon() && fabs(beta) < std::numeric_limits<double>::epsilon();
 
     if( rtype < 0 )
@@ -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());
 
     int sdepth = depth(), ddepth = CV_MAT_DEPTH(rtype);
-    /*if( sdepth == ddepth && noScale )
+    if( sdepth == ddepth && noScale )
     {
         copyTo(dst);
         return;
-    }*/
+    }
 
     GpuMat temp;
     const GpuMat* psrc = this;

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();
+                //TickMeter tm;
+                //tm.start();
+                //for(int i = 0; i < 50; ++i)
                     cpumatsrc.convertTo(cpumatdst, dst_type, alpha, beta);
-                //cput = ((double)getTickCount() - cput) / getTickFrequency();
+                //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();
+                //tm.reset();
+
+                try
+                {
+                    //tm.start();
+                    //for(int i = 0; i < 50; ++i)
                         gpumatsrc.convertTo(gpumatdst, dst_type, alpha, beta);
-                //gput = ((double)getTickCount() - gput) / getTickFrequency();
+                    //tm.stop();
+                    //cout << "\tGPU FPS = " << 50.0/tm.getTimeSec() << endl;
+                }
+                catch(cv::Exception& e)
+                {
+                    cout << "ERROR: " << e.err << endl;
+                    passed = false;
+                    break;
+                }
 
-                /*cout << "convertTo time: " << endl;
+                double r = norm(cpumatdst, gpumatdst, NORM_INF);
-                cout << "CPU time: " << cput << endl;
-                cout << "GPU time: " << gput << endl;/**/               
-
-                double r = norm(cpumatdst, gpumatdst, NORM_L1);
                 if (r > 1)
                 {                    
-                    /*namedWindow("CPU"); 
+                    cout << "FAILED: " << "SRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << " NORM = " << r << endl;
-                    imshow("CPU", cpumatdst); 
-                    namedWindow("GPU"); 
-                    imshow("GPU", gpumatdst); 
-                    waitKey();/**/
-                    
-                    cout << "Failed:" << endl;
-                    cout << "\tr = " << r << endl;
-                    cout << "\tSRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << endl;/**/
 
                     passed = false;
                 }