From 5285722c1c2c1d15bbfc132f3c202449aaea6b49 Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 27 Sep 2010 14:10:19 +0000 Subject: [PATCH] added gpu::magnitude for complex source vector (two channels float). added gpu::cvtColor for BGR5x5 <-> BGR and BGR5x5 <-> Gray. --- modules/gpu/include/opencv2/gpu/gpu.hpp | 2 + modules/gpu/src/arithm.cpp | 22 ++- modules/gpu/src/cuda/color.cu | 272 +++++++++++++++++--------------- modules/gpu/src/imgproc_gpu.cpp | 53 ++++--- tests/gpu/src/gputest_main.cpp | 6 +- tests/gpu/src/imgproc_gpu.cpp | 12 +- 6 files changed, 202 insertions(+), 165 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 51abde7..314a8c2 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -452,6 +452,8 @@ namespace cv //! computes magnitude (magnitude(i)) of each (x(i), y(i)) vector CV_EXPORTS void magnitude(const GpuMat& x, const GpuMat& y, GpuMat& magnitude); + //! computes magnitude (magnitude(i)) of complex (x(i).re, x(i).im) vector + CV_EXPORTS void magnitude(const GpuMat& x, GpuMat& magnitude); ////////////////////////////// Image processing ////////////////////////////// diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp index 7f99912..00e08a1 100644 --- a/modules/gpu/src/arithm.cpp +++ b/modules/gpu/src/arithm.cpp @@ -70,6 +70,7 @@ void cv::gpu::LUT(const GpuMat&, const Mat&, GpuMat&) { throw_nogpu(); } void cv::gpu::exp(const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::log(const GpuMat&, GpuMat&) { throw_nogpu(); } void cv::gpu::magnitude(const GpuMat&, const GpuMat&, GpuMat&) { throw_nogpu(); } +void cv::gpu::magnitude(const GpuMat&, GpuMat&) { throw_nogpu(); } #else /* !defined (HAVE_CUDA) */ @@ -530,6 +531,19 @@ void cv::gpu::log(const GpuMat& src, GpuMat& dst) //////////////////////////////////////////////////////////////////////// // magnitude +void cv::gpu::magnitude(const GpuMat& src, GpuMat& dst) +{ + CV_Assert(src.type() == CV_32FC2); + + dst.create(src.size(), CV_32FC1); + + NppiSize sz; + sz.width = src.cols; + sz.height = src.rows; + + nppSafeCall( nppiMagnitude_32fc32f_C1R(src.ptr(), src.step, dst.ptr(), dst.step, sz) ); +} + void cv::gpu::magnitude(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) { CV_DbgAssert(src1.type() == src2.type() && src1.size() == src2.size()); @@ -539,13 +553,7 @@ void cv::gpu::magnitude(const GpuMat& src1, const GpuMat& src2, GpuMat& dst) GpuMat srcs[] = {src1, src2}; cv::gpu::merge(srcs, 2, src); - dst.create(src1.size(), src1.type()); - - NppiSize sz; - sz.width = src.cols; - sz.height = src.rows; - - nppSafeCall( nppiMagnitude_32fc32f_C1R(src.ptr(), src.step, dst.ptr(), dst.step, sz) ); + cv::gpu::magnitude(src, dst); } #endif /* !defined (HAVE_CUDA) */ \ No newline at end of file diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu index 418518e..dadb959 100644 --- a/modules/gpu/src/cuda/color.cu +++ b/modules/gpu/src/cuda/color.cu @@ -312,41 +312,6 @@ namespace imgproc } } - /*struct RGB5x52RGB - { - typedef uchar channel_type; - - RGB5x52RGB(int _dstcn, int _blueIdx, int _greenBits) - : dstcn(_dstcn), blueIdx(_blueIdx), greenBits(_greenBits) {} - - void operator()(const uchar* src, uchar* dst, int n) const - { - int dcn = dstcn, bidx = blueIdx; - if( greenBits == 6 ) - for( int i = 0; i < n; i++, dst += dcn ) - { - unsigned t = ((const unsigned short*)src)[i]; - dst[bidx] = (uchar)(t << 3); - dst[1] = (uchar)((t >> 3) & ~3); - dst[bidx ^ 2] = (uchar)((t >> 8) & ~7); - if( dcn == 4 ) - dst[3] = 255; - } - else - for( int i = 0; i < n; i++, dst += dcn ) - { - unsigned t = ((const unsigned short*)src)[i]; - dst[bidx] = (uchar)(t << 3); - dst[1] = (uchar)((t >> 2) & ~7); - dst[bidx ^ 2] = (uchar)((t >> 7) & ~7); - if( dcn == 4 ) - dst[3] = t & 0x8000 ? 255 : 0; - } - } - - int dstcn, blueIdx, greenBits; - };*/ - template struct RGB2RGB5x5Converter {}; template struct RGB2RGB5x5Converter @@ -415,7 +380,7 @@ namespace cv { namespace gpu { namespace improc {RGB5x52RGB_caller<6, 3>, RGB5x52RGB_caller<6, 4>} }; - RGB5x52RGB_callers[green_bits - 5][dstcn - 5](src, dst, bidx, stream); + RGB5x52RGB_callers[green_bits - 5][dstcn - 3](src, dst, bidx, stream); } template @@ -471,28 +436,37 @@ namespace imgproc } } - //struct Gray2RGB5x5 - //{ - // typedef uchar channel_type; - // - // Gray2RGB5x5(int _greenBits) : greenBits(_greenBits) {} - // void operator()(const uchar* src, uchar* dst, int n) const - // { - // if( greenBits == 6 ) - // for( int i = 0; i < n; i++ ) - // { - // int t = src[i]; - // ((unsigned short*)dst)[i] = (unsigned short)((t >> 3)|((t & ~3) << 3)|((t & ~7) << 8)); - // } - // else - // for( int i = 0; i < n; i++ ) - // { - // int t = src[i] >> 3; - // ((unsigned short*)dst)[i] = (unsigned short)(t|(t << 5)|(t << 10)); - // } - // } - // int greenBits; - //}; + template struct Gray2RGB5x5Converter {}; + + template<> struct Gray2RGB5x5Converter<6> + { + static __device__ unsigned short cvt(unsigned int t) + { + return (unsigned short)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8)); + } + }; + template<> struct Gray2RGB5x5Converter<5> + { + static __device__ unsigned short cvt(unsigned int t) + { + t >>= 3; + return (unsigned short)(t | (t << 5) | (t << 10)); + } + }; + + template + __global__ void Gray2RGB5x5(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols) + { + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (y < rows && x < cols) + { + unsigned int src = src_[y * src_step + x]; + + *(unsigned short*)(dst_ + y * dst_step + (x << 1)) = Gray2RGB5x5Converter::cvt(src); + } + } } namespace cv { namespace gpu { namespace improc @@ -536,60 +510,86 @@ namespace cv { namespace gpu { namespace improc Gray2RGB_callers[dstcn - 3](src, dst, stream); } + + template + void Gray2RGB5x5_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream) + { + dim3 threads(32, 8, 1); + dim3 grid(1, 1, 1); + + grid.x = divUp(src.cols, threads.x); + grid.y = divUp(src.rows, threads.y); + + imgproc::Gray2RGB5x5<<>>(src.ptr, src.step, + dst.ptr, dst.step, src.rows, src.cols); + + if (stream == 0) + cudaSafeCall( cudaThreadSynchronize() ); + } + + void Gray2RGB5x5_gpu(const DevMem2D& src, const DevMem2D& dst, int green_bits, cudaStream_t stream) + { + typedef void (*Gray2RGB5x5_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream); + static const Gray2RGB5x5_caller_t Gray2RGB5x5_callers[2] = + { + Gray2RGB5x5_caller<5>, Gray2RGB5x5_caller<6> + }; + + Gray2RGB5x5_callers[green_bits - 5](src, dst, stream); + } }}} ///////////////////////////////// Color to Grayscale //////////////////////////////// namespace imgproc { - //#undef R2Y - //#undef G2Y - //#undef B2Y - // - //enum - //{ - // yuv_shift = 14, - // xyz_shift = 12, - // R2Y = 4899, - // G2Y = 9617, - // B2Y = 1868, - // BLOCK_SIZE = 256 - //}; - - //struct RGB5x52Gray - //{ - // typedef uchar channel_type; - // - // RGB5x52Gray(int _greenBits) : greenBits(_greenBits) {} - // void operator()(const uchar* src, uchar* dst, int n) const - // { - // if( greenBits == 6 ) - // for( int i = 0; i < n; i++ ) - // { - // int t = ((unsigned short*)src)[i]; - // dst[i] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + - // ((t >> 3) & 0xfc)*G2Y + - // ((t >> 8) & 0xf8)*R2Y, yuv_shift); - // } - // else - // for( int i = 0; i < n; i++ ) - // { - // int t = ((unsigned short*)src)[i]; - // dst[i] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y + - // ((t >> 2) & 0xf8)*G2Y + - // ((t >> 7) & 0xf8)*R2Y, yuv_shift); - // } - // } - // int greenBits; - //}; + #undef R2Y + #undef G2Y + #undef B2Y + + enum + { + yuv_shift = 14, + xyz_shift = 12, + R2Y = 4899, + G2Y = 9617, + B2Y = 1868, + BLOCK_SIZE = 256 + }; - __global__ void RGB2Gray_3(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx) + template struct RGB5x52GrayConverter {}; + + template<> struct RGB5x52GrayConverter<6> + { + static __device__ unsigned char cvt(unsigned int t) + { + return (unsigned char)CV_DESCALE(((t << 3) & 0xf8)*B2Y + ((t >> 3) & 0xfc)*G2Y + ((t >> 8) & 0xf8)*R2Y, yuv_shift); + } + }; + template<> struct RGB5x52GrayConverter<5> + { + static __device__ unsigned char cvt(unsigned int t) + { + return (unsigned char)CV_DESCALE(((t << 3) & 0xf8)*B2Y + ((t >> 2) & 0xf8)*G2Y + ((t >> 7) & 0xf8)*R2Y, yuv_shift); + } + }; + + template + __global__ void RGB5x52Gray(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols) { - const int cr = 4899; - const int cg = 9617; - const int cb = 1868; - const int yuv_shift = 14; + const int x = blockDim.x * blockIdx.x + threadIdx.x; + const int y = blockDim.y * blockIdx.y + threadIdx.y; + + if (y < rows && x < cols) + { + unsigned int src = *(unsigned short*)(src_ + y * src_step + (x << 1)); + dst_[y * dst_step + x] = RGB5x52GrayConverter::cvt(src); + } + } + + __global__ void RGB2Gray_3(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx) + { const int x = (blockDim.x * blockIdx.x + threadIdx.x) << 2; const int y = blockDim.y * blockIdx.y + threadIdx.y; @@ -600,19 +600,19 @@ namespace imgproc uchar t0 = src[bidx], t1 = src[1], t2 = src[bidx ^ 2]; uchar4 dst; - dst.x = (uchar)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.x = (uchar)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); src += 3; t0 = src[bidx], t1 = src[1], t2 = src[bidx ^ 2]; - dst.y = (uchar)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.y = (uchar)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); src += 3; t0 = src[bidx], t1 = src[1], t2 = src[bidx ^ 2]; - dst.z = (uchar)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.z = (uchar)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); src += 3; t0 = src[bidx], t1 = src[1], t2 = src[bidx ^ 2]; - dst.w = (uchar)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.w = (uchar)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); *(uchar4*)(dst_ + y * dst_step + x) = dst; } @@ -620,11 +620,6 @@ namespace imgproc __global__ void RGB2Gray_3(const unsigned short* src_, size_t src_step, unsigned short* dst_, size_t dst_step, int rows, int cols, int bidx) { - const int cr = 4899; - const int cg = 9617; - const int cb = 1868; - const int yuv_shift = 14; - const int x = (blockDim.x * blockIdx.x + threadIdx.x) << 1; const int y = blockDim.y * blockIdx.y + threadIdx.y; @@ -635,11 +630,11 @@ namespace imgproc unsigned short t0 = src[bidx], t1 = src[1], t2 = src[bidx ^ 2]; ushort2 dst; - dst.x = (unsigned short)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.x = (unsigned short)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); src += 3; t0 = src[bidx], t1 = src[1], t2 = src[bidx ^ 2]; - dst.y = (unsigned short)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.y = (unsigned short)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); *(ushort2*)(dst_ + y * dst_step + x) = dst; } @@ -665,11 +660,6 @@ namespace imgproc __global__ void RGB2Gray_4(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx) { - const int cr = 4899; - const int cg = 9617; - const int cb = 1868; - const int yuv_shift = 14; - const int x = (blockDim.x * blockIdx.x + threadIdx.x) << 2; const int y = blockDim.y * blockIdx.y + threadIdx.y; @@ -680,19 +670,19 @@ namespace imgproc uchar t0 = ((uchar*)(&src))[bidx], t1 = src.y, t2 = ((uchar*)(&src))[bidx ^ 2]; uchar4 dst; - dst.x = (uchar)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.x = (uchar)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); src = *(uchar4*)(src_ + y * src_step + (x << 2) + 4); t0 = ((uchar*)(&src))[bidx], t1 = src.y, t2 = ((uchar*)(&src))[bidx ^ 2]; - dst.y = (uchar)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.y = (uchar)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); src = *(uchar4*)(src_ + y * src_step + (x << 2) + 8); t0 = ((uchar*)(&src))[bidx], t1 = src.y, t2 = ((uchar*)(&src))[bidx ^ 2]; - dst.z = (uchar)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.z = (uchar)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); src = *(uchar4*)(src_ + y * src_step + (x << 2) + 12); t0 = ((uchar*)(&src))[bidx], t1 = src.y, t2 = ((uchar*)(&src))[bidx ^ 2]; - dst.w = (uchar)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.w = (uchar)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); *(uchar4*)(dst_ + y * dst_step + x) = dst; } @@ -700,11 +690,6 @@ namespace imgproc __global__ void RGB2Gray_4(const unsigned short* src_, size_t src_step, unsigned short* dst_, size_t dst_step, int rows, int cols, int bidx) { - const int cr = 4899; - const int cg = 9617; - const int cb = 1868; - const int yuv_shift = 14; - const int x = (blockDim.x * blockIdx.x + threadIdx.x) << 1; const int y = blockDim.y * blockIdx.y + threadIdx.y; @@ -715,11 +700,11 @@ namespace imgproc unsigned short t0 = ((unsigned short*)(&src))[bidx], t1 = src.y, t2 = ((unsigned short*)(&src))[bidx ^ 2]; ushort2 dst; - dst.x = (unsigned short)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.x = (unsigned short)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); src = *(ushort4*)(src_ + y * src_step + (x << 2) + 4); t0 = ((unsigned short*)(&src))[bidx], t1 = src.y, t2 = ((unsigned short*)(&src))[bidx ^ 2]; - dst.y = (unsigned short)CV_DESCALE((unsigned)(t0 * cb + t1 * cg + t2 * cr), yuv_shift); + dst.y = (unsigned short)CV_DESCALE((unsigned)(t0 * B2Y + t1 * G2Y + t2 * R2Y), yuv_shift); *(ushort2*)(dst_ + y * dst_step + x) = dst; } @@ -820,6 +805,33 @@ namespace cv { namespace gpu { namespace improc if (stream == 0) cudaSafeCall( cudaThreadSynchronize() ); } + + template + void RGB5x52Gray_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream) + { + dim3 threads(32, 8, 1); + dim3 grid(1, 1, 1); + + grid.x = divUp(src.cols, threads.x); + grid.y = divUp(src.rows, threads.y); + + imgproc::RGB5x52Gray<<>>(src.ptr, src.step, + dst.ptr, dst.step, src.rows, src.cols); + + if (stream == 0) + cudaSafeCall( cudaThreadSynchronize() ); + } + + void RGB5x52Gray_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, cudaStream_t stream) + { + typedef void (*RGB5x52Gray_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream); + static const RGB5x52Gray_caller_t RGB5x52Gray_callers[2] = + { + RGB5x52Gray_caller<5>, RGB5x52Gray_caller<6> + }; + + RGB5x52Gray_callers[green_bits - 5](src, dst, stream); + } }}} ///////////////////////////////////// RGB <-> YCrCb ////////////////////////////////////// diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp index 750044a..d7c47da 100644 --- a/modules/gpu/src/imgproc_gpu.cpp +++ b/modules/gpu/src/imgproc_gpu.cpp @@ -95,10 +95,12 @@ namespace cv { namespace gpu void Gray2RGB_gpu(const DevMem2D& src, const DevMem2D& dst, int dstcn, cudaStream_t stream); void Gray2RGB_gpu(const DevMem2D_& src, const DevMem2D_& dst, int dstcn, cudaStream_t stream); void Gray2RGB_gpu(const DevMem2Df& src, const DevMem2Df& dst, int dstcn, cudaStream_t stream); + void Gray2RGB5x5_gpu(const DevMem2D& src, const DevMem2D& dst, int green_bits, cudaStream_t stream); void RGB2Gray_gpu(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream); void RGB2Gray_gpu(const DevMem2D_& src, int srccn, const DevMem2D_& dst, int bidx, cudaStream_t stream); void RGB2Gray_gpu(const DevMem2Df& src, int srccn, const DevMem2Df& dst, int bidx, cudaStream_t stream); + void RGB5x52Gray_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, cudaStream_t stream); } }} @@ -267,18 +269,20 @@ namespace stream); break; - //case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: - //case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: - // if(dcn <= 0) dcn = 3; - // CV_Assert( (dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U ); - // out.create(sz, CV_MAKETYPE(depth, dcn)); - - // improc::RGB5x52RGB_gpu(src, code == CV_BGR2BGR565 || code == CV_RGB2BGR565 || - // code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5, out, dcn, - // code == CV_BGR2BGR565 || code == CV_BGR2BGR555 || - // code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2, - // stream); - // break; + case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: + case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: + if(dcn <= 0) dcn = 3; + CV_Assert( (dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U ); + out.create(sz, CV_MAKETYPE(depth, dcn)); + + improc::RGB5x52RGB_gpu(src, + code == CV_BGR5652BGR || code == CV_BGR5652RGB || + code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5, + out, dcn, + code == CV_BGR5652BGR || code == CV_BGR5552BGR || + code == CV_BGR5652BGRA || code == CV_BGR5552BGRA ? 0 : 2, + stream); + break; case CV_BGR2GRAY: case CV_BGRA2GRAY: case CV_RGB2GRAY: case CV_RGBA2GRAY: CV_Assert(scn == 3 || scn == 4); @@ -294,11 +298,13 @@ namespace improc::RGB2Gray_gpu((DevMem2Df)src, scn, (DevMem2Df)out, bidx, stream); break; - //case CV_BGR5652GRAY: case CV_BGR5552GRAY: - // CV_Assert( scn == 2 && depth == CV_8U ); - // dst.create(sz, CV_8UC1); - // CvtColorLoop(src, dst, RGB5x52Gray(code == CV_BGR5652GRAY ? 6 : 5)); - // break; + case CV_BGR5652GRAY: case CV_BGR5552GRAY: + CV_Assert( scn == 2 && depth == CV_8U ); + + out.create(sz, CV_8UC1); + + improc::RGB5x52Gray_gpu(src, code == CV_BGR5652GRAY ? 6 : 5, out, stream); + break; case CV_GRAY2BGR: case CV_GRAY2BGRA: if (dcn <= 0) @@ -315,12 +321,13 @@ namespace improc::Gray2RGB_gpu((DevMem2Df)src, (DevMem2Df)out, dcn, stream); break; - //case CV_GRAY2BGR565: case CV_GRAY2BGR555: - // CV_Assert( scn == 1 && depth == CV_8U ); - // dst.create(sz, CV_8UC2); - // - // CvtColorLoop(src, dst, Gray2RGB5x5(code == CV_GRAY2BGR565 ? 6 : 5)); - // break; + case CV_GRAY2BGR565: case CV_GRAY2BGR555: + CV_Assert( scn == 1 && depth == CV_8U ); + + out.create(sz, CV_8UC2); + + improc::Gray2RGB5x5_gpu(src, out, code == CV_GRAY2BGR565 ? 6 : 5, stream); + break; case CV_RGB2YCrCb: CV_Assert(scn == 3 && depth == CV_8U); diff --git a/tests/gpu/src/gputest_main.cpp b/tests/gpu/src/gputest_main.cpp index d4b9b3f..463b2b7 100644 --- a/tests/gpu/src/gputest_main.cpp +++ b/tests/gpu/src/gputest_main.cpp @@ -47,9 +47,9 @@ const char* blacklist[] = { "GPU-NppImageSum", // crash "GPU-MatOperatorAsyncCall", // crash - //"GPU-NppErode", // npp func returns error code (CUDA_KERNEL_LAUNCH_ERROR or TEXTURE_BIND_ERROR) - //"GPU-NppDilate", // npp func returns error code (CUDA_KERNEL_LAUNCH_ERROR or TEXTURE_BIND_ERROR) - //"GPU-NppMorphologyEx", // npp func returns error code (CUDA_KERNEL_LAUNCH_ERROR or TEXTURE_BIND_ERROR) + "GPU-NppErode", // npp func returns error code (CUDA_KERNEL_LAUNCH_ERROR or TEXTURE_BIND_ERROR) + "GPU-NppDilate", // npp func returns error code (CUDA_KERNEL_LAUNCH_ERROR or TEXTURE_BIND_ERROR) + "GPU-NppMorphologyEx", // npp func returns error code (CUDA_KERNEL_LAUNCH_ERROR or TEXTURE_BIND_ERROR) //"GPU-NppImageDivide", // different round mode //"GPU-NppImageMeanStdDev", // different precision //"GPU-NppImageMinNax", // npp bug diff --git a/tests/gpu/src/imgproc_gpu.cpp b/tests/gpu/src/imgproc_gpu.cpp index 0f8a6ae..beff852 100644 --- a/tests/gpu/src/imgproc_gpu.cpp +++ b/tests/gpu/src/imgproc_gpu.cpp @@ -498,8 +498,16 @@ void CV_GpuCvtColorTest::run( int ) try { //run tests - int codes[] = { CV_BGR2RGB, /* CV_RGB2YCrCb, CV_YCrCb2RGB,*/ CV_RGB2RGBA, CV_RGBA2BGRA, CV_BGRA2GRAY, CV_GRAY2RGB, CV_RGB2BGR555/*, CV_BGR5552BGR/*, CV_BGR2BGR565, CV_BGR5652RGB*/}; - const char* codes_str[] = {"CV_BGR2RGB", /*"CV_RGB2YCrCb", "CV_YCrCb2RGB",*/ "CV_RGB2RGBA", "CV_RGBA2BGRA", "CV_BGRA2GRAY", "CV_GRAY2RGB", "CV_RGB2BGR555"/*, "CV_BGR5552BGR"/*, "CV_BGR2BGR565", "CV_BGR5652RGB"*/}; + int codes[] = { CV_BGR2RGB, CV_RGB2BGRA, CV_BGRA2RGB, + CV_RGB2BGR555, CV_BGR5552BGR, CV_BGR2BGR565, CV_BGR5652RGB, + /* CV_RGB2YCrCb, CV_YCrCb2RGB,*/ + CV_RGB2GRAY, CV_GRAY2BGRA, CV_BGRA2GRAY, + CV_GRAY2BGR555, CV_BGR5552GRAY, CV_GRAY2BGR565, CV_BGR5652GRAY}; + const char* codes_str[] = { "CV_BGR2RGB", "CV_RGB2BGRA", "CV_BGRA2RGB", + "CV_RGB2BGR555", "CV_BGR5552BGR", "CV_BGR2BGR565", "CV_BGR5652RGB", + /* "CV_RGB2YCrCb", "CV_YCrCb2RGB",*/ + "CV_RGB2GRAY", "CV_GRAY2BGRA", "CV_BGRA2GRAY", + "CV_GRAY2BGR555", "CV_BGR5552GRAY", "CV_GRAY2BGR565", "CV_BGR5652GRAY"}; int codes_num = sizeof(codes) / sizeof(int); for (int i = 0; i < codes_num; ++i) -- 2.7.4