From 26c485963451163f6cda6e6481b66f78afc8ea70 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Thu, 22 Jul 2010 14:50:31 +0000 Subject: [PATCH] reduced code convert_to by using templates, merged with copyTo --- modules/gpu/src/cuda/matrix_operations.cu | 429 ++++++++-------------- modules/gpu/src/matrix_operations.cpp | 6 +- tests/gpu/src/convert_to.cpp | 51 +-- 3 files changed, 190 insertions(+), 296 deletions(-) diff --git a/modules/gpu/src/cuda/matrix_operations.cu b/modules/gpu/src/cuda/matrix_operations.cu index ad2e0f3317..df68b997cd 100644 --- 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 __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 __global__ void kernel_set_to_without_mask(T * mat, int cols, int rows, int step, int channels) @@ -103,236 +104,134 @@ namespace mat_operators } - ////////////////////////////////////////////////////////// - // ConvertTo - ////////////////////////////////////////////////////////// +/////////////////////////////////////////////////////////////////////////// +//////////////////////////////// ConvertTo //////////////////////////////// +/////////////////////////////////////////////////////////////////////////// + + template + struct CalcTraits + { + __device__ static DT calc(T src, double alpha, double beta) + { + return (DT)__double2int_rn(alpha * src + beta); + } + }; + template + struct CalcTraits + { + __device__ static float calc(T src, double alpha, double beta) + { + return (float)(alpha * src + beta); + } + }; + template + struct CalcTraits + { + __device__ static double calc(T src, double alpha, double beta) + { + return alpha * src + beta; + } + }; template + struct ConverterTraits + { + enum {shift=1}; + + typedef T read_type; + typedef DT write_type; + }; + template + struct ConverterTraits + { + enum {shift=4}; + + typedef char4 read_type; + typedef char4 write_type; + }; + template + struct ConverterTraits + { + enum {shift=4}; + + typedef short4 read_type; + typedef char4 write_type; + }; + template + struct ConverterTraits + { + enum {shift=4}; + + typedef int4 read_type; + typedef char4 write_type; + }; + template + struct ConverterTraits + { + enum {shift=2}; + + typedef char2 read_type; + typedef short2 write_type; + }; + template + struct ConverterTraits + { + enum {shift=2}; + + typedef short2 read_type; + typedef short2 write_type; + }; + template + struct ConverterTraits + { + enum {shift=2}; + + typedef int2 read_type; + typedef short2 write_type; + }; + + template 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) + if (y < height) { const T* src = (const T*)(srcmat + src_step * y); DT* dst = (DT*)(dstmat + dst_step * y); + if ((x * ConverterTraits::shift) + ConverterTraits::shift - 1 < width) + { + typename ConverterTraits::read_type srcn_el = ((const typename ConverterTraits::read_type*)src)[x]; + typename ConverterTraits::write_type dstn_el; - dst[x] = (DT)__double2int_rn(alpha * src[x] + beta); + const T* src1_el = (const T*) &srcn_el; + DT* dst1_el = (DT*) &dstn_el; + + for (int i = 0; i < ConverterTraits::shift; ++i) + dst1_el[i] = CalcTraits::calc(src1_el[i], alpha, beta); + + ((typename ConverterTraits::write_type*)dst)[x] = dstn_el; + } + else + { + for (int i = 0; i < ConverterTraits::shift - 1; ++i) + if ((x * ConverterTraits::shift) + i < width) + dst[(x * ConverterTraits::shift) + i] = CalcTraits::calc(src[(x * ConverterTraits::shift) + i], alpha, 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::shift), divUp(height, block.y)); } }; - - template - 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 (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); - - 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 - 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 (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 - 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 (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 - 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 (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 - 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); - 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 - 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); - 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 + + template __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::convert(srcmat, src_step, dstmat, dst_step, width, height, alpha, beta); + Converter::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 - ////////////////////////////////////////////////////////////// + +/////////////////////////////////////////////////////////////////////////// +//////////////////////////////// ConvertTo //////////////////////////////// +/////////////////////////////////////////////////////////////////////////// - typedef void (*CvtFunc)(const DevMem2D& src, DevMem2D& dst, size_t width, size_t height, double alpha, double beta); + 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 + 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::calcGrid(width, height, block); + ::mat_operators::kernel_convert_to<<>>(src.ptr, src.step, dst.ptr, dst.step, width, height, alpha, beta); + cudaSafeCall( cudaThreadSynchronize() ); + } - template - 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::calcGrid(width, height, block); - ::mat_operators::kernel_convert_to<<>>(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_, cvt_, cvt_, cvt_, + cvt_, cvt_, cvt_, 0}, - 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_, cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, 0}, + {cvt_, cvt_, cvt_, cvt_, + cvt_, cvt_, cvt_, 0}, - {cvt_, cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, 0}, + {cvt_, cvt_, cvt_, cvt_, + cvt_, cvt_, cvt_, 0}, - {cvt_, cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, 0}, + {cvt_, cvt_, cvt_, cvt_, + cvt_, cvt_, cvt_, 0}, - {cvt_, cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, 0}, + {cvt_, cvt_, cvt_, + cvt_, cvt_, cvt_, cvt_, 0}, - {cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, cvt_, 0}, + {cvt_, cvt_, cvt_, + cvt_, cvt_, cvt_, cvt_, 0}, - {cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, cvt_, 0}, + {cvt_, cvt_, cvt_, + cvt_, cvt_, cvt_, cvt_, 0}, - {cvt_, cvt_, cvt_, - cvt_, cvt_, cvt_, cvt_, 0}, + {0,0,0,0,0,0,0,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); - } - } - - - } -} + 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 9469f59185..f64733d613 100644 --- 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 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::epsilon() && fabs(beta) < std::numeric_limits::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; diff --git a/tests/gpu/src/convert_to.cpp b/tests/gpu/src/convert_to.cpp index c53fbc455a..b2e231a065 100644 --- 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;/**/ - - double r = norm(cpumatdst, gpumatdst, NORM_L1); - if (r > 1) + try { - /*namedWindow("CPU"); - 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;/**/ + //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_INF); + if (r > 1) + { + cout << "FAILED: " << "SRC_TYPE=" << types_str[i] << "C" << c << " DST_TYPE=" << types_str[j] << " NORM = " << r << endl; passed = false; }