 
         typedef char4 read_type;
         typedef char4 write_type;
-    };    
+    };
     template <typename T, typename DT>
     struct ReadWriteTraits<T, DT, 2, 1>
     {

        typedef int2 read_type;
        typedef short2 write_type;
    };
    
    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)
    {
        typedef typename ReadWriteTraits<T, DT, sizeof(T), sizeof(DT)>::read_type read_type;


                            SetToFunc_with_mask func = tab[depth];

                            if (func == 0) cv::gpu::error("Unsupported convert operation", __FILE__, __LINE__);

                            func(mat, mask, channels);
                        }


///////////////////////////////////////////////////////////////////////////
//////////////////////////////// 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)
				    {
					const int shift = ::mat_operators::ReadWriteTraits<T, DT, sizeof(T), sizeof(DT)>::shift;

                                        dim3 block(32, 8);
                                        dim3 grid(divUp(width, block.x * shift), divUp(height, block.y));

                                        ::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("Unsupported convert operation", __FILE__, __LINE__);
					    func(src, dst, width, height, alpha, beta);
					}
				} // namespace impl
	    } // namespace gpu
    } // namespace cv
