From 051adcb78606effd24fdcb9842d19c568c67084e Mon Sep 17 00:00:00 2001 From: Vladislav Vinogradov Date: Mon, 30 Jul 2012 15:18:48 +0400 Subject: [PATCH] added gpu BGR<->Lab and RGB<->Luv color conversion and gammaCorrection --- modules/gpu/include/opencv2/gpu/gpu.hpp | 49 ++++---- modules/gpu/src/color.cpp | 160 +++++++++++++++++++++++- modules/gpu/test/test_color.cpp | 46 +++++++ 3 files changed, 228 insertions(+), 27 deletions(-) diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp index 515a4a2751..f6d869435c 100644 --- a/modules/gpu/include/opencv2/gpu/gpu.hpp +++ b/modules/gpu/include/opencv2/gpu/gpu.hpp @@ -622,6 +622,9 @@ CV_EXPORTS void cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn = 0, //! channel order. CV_EXPORTS void swapChannels(GpuMat& image, const int dstOrder[4], Stream& stream = Stream::Null()); +//! Routines for correcting image color gamma +CV_EXPORTS void gammaCorrection(const GpuMat& src, GpuMat& dst, bool forward = true, Stream& stream = Stream::Null()); + //! applies fixed threshold to the image CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh, double maxval, int type, Stream& stream = Stream::Null()); @@ -1411,7 +1414,7 @@ public: }; ////////////////////////////////// CascadeClassifier_GPU ////////////////////////////////////////// -// The cascade classifier class for object detection: supports old haar and new lbp xlm formats and nvbin for haar cascades olny. +// The cascade classifier class for object detection: supports old haar and new lbp xlm formats and nvbin for haar cascades olny. class CV_EXPORTS CascadeClassifier_GPU { public: @@ -1421,28 +1424,28 @@ public: bool empty() const; bool load(const std::string& filename); - void release(); - - /* returns number of detected objects */ - int detectMultiScale(const GpuMat& image, GpuMat& objectsBuf, double scaleFactor = 1.2, int minNeighbors = 4, Size minSize = Size()); - - bool findLargestObject; - bool visualizeInPlace; - - Size getClassifierSize() const; - -private: - struct CascadeClassifierImpl; - CascadeClassifierImpl* impl; - struct HaarCascade; - struct LbpCascade; - friend class CascadeClassifier_GPU_LBP; - -public: - int detectMultiScale(const GpuMat& image, GpuMat& objectsBuf, Size maxObjectSize, Size minSize = Size(), double scaleFactor = 1.1, int minNeighbors = 4); -}; - -////////////////////////////////// SURF ////////////////////////////////////////// + void release(); + + /* returns number of detected objects */ + int detectMultiScale(const GpuMat& image, GpuMat& objectsBuf, double scaleFactor = 1.2, int minNeighbors = 4, Size minSize = Size()); + + bool findLargestObject; + bool visualizeInPlace; + + Size getClassifierSize() const; + +private: + struct CascadeClassifierImpl; + CascadeClassifierImpl* impl; + struct HaarCascade; + struct LbpCascade; + friend class CascadeClassifier_GPU_LBP; + +public: + int detectMultiScale(const GpuMat& image, GpuMat& objectsBuf, Size maxObjectSize, Size minSize = Size(), double scaleFactor = 1.1, int minNeighbors = 4); +}; + +////////////////////////////////// SURF ////////////////////////////////////////// class CV_EXPORTS SURF_GPU { diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp index a47758c873..1fe8e87f44 100644 --- a/modules/gpu/src/color.cpp +++ b/modules/gpu/src/color.cpp @@ -49,6 +49,7 @@ using namespace cv::gpu; void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, Stream&) { throw_nogpu(); } void cv::gpu::swapChannels(GpuMat&, const int[], Stream&) { throw_nogpu(); } +void cv::gpu::gammaCorrection(const GpuMat&, GpuMat&, bool, Stream&) { throw_nogpu(); } #else /* !defined (HAVE_CUDA) */ @@ -1142,6 +1143,116 @@ namespace funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream)); } + + void bgr_to_lab(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + #if (CUDA_VERSION < 5000) + (void)src; + (void)dst; + (void)dcn; + (void)stream; + CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" ); + #else + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 3); + + dcn = src.channels(); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + NppStreamHandler h(StreamAccessor::getStream(stream)); + + NppiSize oSizeROI; + oSizeROI.width = src.cols; + oSizeROI.height = src.rows; + + nppSafeCall( nppiBGRToLab_8u_C3R(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), oSizeROI) ); + #endif + } + + void lab_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + #if (CUDA_VERSION < 5000) + (void)src; + (void)dst; + (void)dcn; + (void)stream; + CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" ); + #else + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 3); + + dcn = src.channels(); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + NppStreamHandler h(StreamAccessor::getStream(stream)); + + NppiSize oSizeROI; + oSizeROI.width = src.cols; + oSizeROI.height = src.rows; + + nppSafeCall( nppiLabToBGR_8u_C3R(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), oSizeROI) ); + #endif + } + + void rgb_to_luv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + #if (CUDA_VERSION < 5000) + (void)src; + (void)dst; + (void)dcn; + (void)stream; + CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" ); + #else + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 3 || src.channels() == 4); + + dcn = src.channels(); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + NppStreamHandler h(StreamAccessor::getStream(stream)); + + NppiSize oSizeROI; + oSizeROI.width = src.cols; + oSizeROI.height = src.rows; + + if (dcn == 3) + nppSafeCall( nppiRGBToLUV_8u_C3R(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), oSizeROI) ); + else + nppSafeCall( nppiRGBToLUV_8u_AC4R(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), oSizeROI) ); + #endif + } + + void luv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream) + { + #if (CUDA_VERSION < 5000) + (void)src; + (void)dst; + (void)dcn; + (void)stream; + CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" ); + #else + CV_Assert(src.depth() == CV_8U); + CV_Assert(src.channels() == 3 || src.channels() == 4); + + dcn = src.channels(); + + dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn)); + + NppStreamHandler h(StreamAccessor::getStream(stream)); + + NppiSize oSizeROI; + oSizeROI.width = src.cols; + oSizeROI.height = src.rows; + + if (dcn == 3) + nppSafeCall( nppiLUVToRGB_8u_C3R(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), oSizeROI) ); + else + nppSafeCall( nppiLUVToRGB_8u_AC4R(src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), oSizeROI) ); + #endif + } } void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream& stream) @@ -1203,7 +1314,7 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream 0, // =42 0, // =43 - 0, // CV_BGR2Lab =44 + bgr_to_lab, // CV_BGR2Lab =44 0, // CV_RGB2Lab =45 0, // CV_BayerBG2BGR =46 @@ -1212,7 +1323,7 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream 0, // CV_BayerGR2BGR =49 0, // CV_BGR2Luv =50 - 0, // CV_RGB2Luv =51 + rgb_to_luv, // CV_RGB2Luv =51 bgr_to_hls, // CV_BGR2HLS =52 rgb_to_hls, // CV_RGB2HLS =53 @@ -1220,10 +1331,10 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream hsv_to_bgr, // CV_HSV2BGR =54 hsv_to_rgb, // CV_HSV2RGB =55 - 0, // CV_Lab2BGR =56 + lab_to_bgr, // CV_Lab2BGR =56 0, // CV_Lab2RGB =57 0, // CV_Luv2BGR =58 - 0, // CV_Luv2RGB =59 + luv_to_rgb, // CV_Luv2RGB =59 hls_to_bgr, // CV_HLS2BGR =60 hls_to_rgb, // CV_HLS2RGB =61 @@ -1292,4 +1403,45 @@ void cv::gpu::swapChannels(GpuMat& image, const int dstOrder[4], Stream& s) cudaSafeCall( cudaDeviceSynchronize() ); } +void cv::gpu::gammaCorrection(const GpuMat& src, GpuMat& dst, bool forward, Stream& stream) +{ +#if (CUDA_VERSION < 5000) + (void)src; + (void)dst; + (void)forward; + (void)stream; + CV_Error( CV_StsNotImplemented, "This function works only with CUDA 5.0 or higher" ); +#else + typedef NppStatus (*func_t)(const Npp8u* pSrc, int nSrcStep, Npp8u* pDst, int nDstStep, NppiSize oSizeROI); + typedef NppStatus (*func_inplace_t)(Npp8u* pSrcDst, int nSrcDstStep, NppiSize oSizeROI); + + static const func_t funcs[2][5] = + { + {0, 0, 0, nppiGammaInv_8u_C3R, nppiGammaInv_8u_AC4R}, + {0, 0, 0, nppiGammaFwd_8u_C3R, nppiGammaFwd_8u_AC4R} + }; + static const func_inplace_t funcs_inplace[2][5] = + { + {0, 0, 0, nppiGammaInv_8u_C3IR, nppiGammaInv_8u_AC4IR}, + {0, 0, 0, nppiGammaFwd_8u_C3IR, nppiGammaFwd_8u_AC4IR} + }; + + CV_Assert(src.type() == CV_8UC3 || src.type() == CV_8UC4); + + dst.create(src.size(), src.type()); + + NppStreamHandler h(StreamAccessor::getStream(stream)); + + NppiSize oSizeROI; + oSizeROI.width = src.cols; + oSizeROI.height = src.rows; + + if (dst.data == src.data) + funcs_inplace[forward][src.channels()](dst.ptr(), static_cast(src.step), oSizeROI); + else + funcs[forward][src.channels()](src.ptr(), static_cast(src.step), dst.ptr(), static_cast(dst.step), oSizeROI); + +#endif +} + #endif /* !defined (HAVE_CUDA) */ diff --git a/modules/gpu/test/test_color.cpp b/modules/gpu/test/test_color.cpp index 1d3ced0220..1c08d6ac66 100644 --- a/modules/gpu/test/test_color.cpp +++ b/modules/gpu/test/test_color.cpp @@ -1609,6 +1609,52 @@ TEST_P(CvtColor, RGBA2YUV4) EXPECT_MAT_NEAR(dst_gold, h_dst, 1e-5); } +TEST_P(CvtColor, BGR2Lab) +{ + if (depth != CV_8U) + return; + + try + { + cv::Mat src = readImage("stereobm/aloe-L.png"); + + cv::gpu::GpuMat dst_lab = createMat(src.size(), src.type(), useRoi); + cv::gpu::cvtColor(loadMat(src, useRoi), dst_lab, cv::COLOR_BGR2Lab); + + cv::gpu::GpuMat dst_bgr = createMat(src.size(), src.type(), useRoi); + cv::gpu::cvtColor(dst_lab, dst_bgr, cv::COLOR_Lab2BGR); + + EXPECT_MAT_NEAR(src, dst_bgr, 10); + } + catch (const cv::Exception& e) + { + ASSERT_EQ(CV_StsBadFlag, e.code); + } +} + +TEST_P(CvtColor, BGR2Luv) +{ + if (depth != CV_8U) + return; + + try + { + cv::Mat src = img; + + cv::gpu::GpuMat dst_luv = createMat(src.size(), src.type(), useRoi); + cv::gpu::cvtColor(loadMat(src, useRoi), dst_luv, cv::COLOR_RGB2Luv); + + cv::gpu::GpuMat dst_rgb = createMat(src.size(), src.type(), useRoi); + cv::gpu::cvtColor(dst_luv, dst_rgb, cv::COLOR_Luv2RGB); + + EXPECT_MAT_NEAR(src, dst_rgb, 10); + } + catch (const cv::Exception& e) + { + ASSERT_EQ(CV_StsBadFlag, e.code); + } +} + INSTANTIATE_TEST_CASE_P(GPU_ImgProc, CvtColor, testing::Combine( ALL_DEVICES, DIFFERENT_SIZES,