added gpu BGR<->Lab and RGB<->Luv color conversion and gammaCorrection

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index 515a4a2..f6d8694 100644 (file)
--- 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.\r
 CV_EXPORTS void swapChannels(GpuMat& image, const int dstOrder[4], Stream& stream = Stream::Null());\r
 \r
+//! Routines for correcting image color gamma\r
+CV_EXPORTS void gammaCorrection(const GpuMat& src, GpuMat& dst, bool forward = true, Stream& stream = Stream::Null());\r
+\r
 //! applies fixed threshold to the image\r
 CV_EXPORTS double threshold(const GpuMat& src, GpuMat& dst, double thresh, double maxval, int type, Stream& stream = Stream::Null());\r
 \r
@@ -1411,7 +1414,7 @@ public:
 };\r
 \r
 ////////////////////////////////// CascadeClassifier_GPU //////////////////////////////////////////\r
-// 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.\r
 class CV_EXPORTS CascadeClassifier_GPU\r
 {\r
 public:\r
@@ -1421,28 +1424,28 @@ public:
 \r
     bool empty() const;\r
     bool load(const std::string& filename);\r
-    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();\r
+\r
+    /* returns number of detected objects */\r
+    int detectMultiScale(const GpuMat& image, GpuMat& objectsBuf, double scaleFactor = 1.2, int minNeighbors = 4, Size minSize = Size());\r
+\r
+    bool findLargestObject;\r
+    bool visualizeInPlace;\r
+\r
+    Size getClassifierSize() const;\r
+\r
+private:\r
+    struct CascadeClassifierImpl;\r
+    CascadeClassifierImpl* impl;\r
+    struct HaarCascade;\r
+    struct LbpCascade;\r
+    friend class CascadeClassifier_GPU_LBP;\r
+\r
+public:\r
+    int detectMultiScale(const GpuMat& image, GpuMat& objectsBuf, Size maxObjectSize, Size minSize = Size(), double scaleFactor = 1.1, int minNeighbors = 4);\r
+};\r
+\r
+////////////////////////////////// SURF //////////////////////////////////////////\r
 \r
 class CV_EXPORTS SURF_GPU\r
 {\r

diff --git a/modules/gpu/src/color.cpp b/modules/gpu/src/color.cpp
index a47758c..1fe8e87 100644 (file)
--- a/modules/gpu/src/color.cpp
+++ b/modules/gpu/src/color.cpp
@@ -49,6 +49,7 @@ using namespace cv::gpu;
 \r
 void cv::gpu::cvtColor(const GpuMat&, GpuMat&, int, int, Stream&) { throw_nogpu(); }\r
 void cv::gpu::swapChannels(GpuMat&, const int[], Stream&) { throw_nogpu(); }\r
+void cv::gpu::gammaCorrection(const GpuMat&, GpuMat&, bool, Stream&) { throw_nogpu(); }\r
 \r
 #else /* !defined (HAVE_CUDA) */\r
 \r
@@ -1142,6 +1143,116 @@ namespace
 \r
         funcs[dcn == 4][src.channels() == 4][src.depth()](src, dst, StreamAccessor::getStream(stream));\r
     }\r
+\r
+    void bgr_to_lab(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
+    {\r
+        #if (CUDA_VERSION < 5000)\r
+            (void)src;\r
+            (void)dst;\r
+            (void)dcn;\r
+            (void)stream;\r
+            CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );\r
+        #else\r
+            CV_Assert(src.depth() == CV_8U);\r
+            CV_Assert(src.channels() == 3);\r
+\r
+            dcn = src.channels();\r
+\r
+            dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn));\r
+\r
+            NppStreamHandler h(StreamAccessor::getStream(stream));\r
+\r
+            NppiSize oSizeROI;\r
+            oSizeROI.width = src.cols;\r
+            oSizeROI.height = src.rows;\r
+\r
+            nppSafeCall( nppiBGRToLab_8u_C3R(src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI) );\r
+        #endif\r
+    }\r
+\r
+    void lab_to_bgr(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
+    {\r
+        #if (CUDA_VERSION < 5000)\r
+            (void)src;\r
+            (void)dst;\r
+            (void)dcn;\r
+            (void)stream;\r
+            CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );\r
+        #else\r
+            CV_Assert(src.depth() == CV_8U);\r
+            CV_Assert(src.channels() == 3);\r
+\r
+            dcn = src.channels();\r
+\r
+            dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn));\r
+\r
+            NppStreamHandler h(StreamAccessor::getStream(stream));\r
+\r
+            NppiSize oSizeROI;\r
+            oSizeROI.width = src.cols;\r
+            oSizeROI.height = src.rows;\r
+\r
+            nppSafeCall( nppiLabToBGR_8u_C3R(src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI) );\r
+        #endif\r
+    }\r
+\r
+    void rgb_to_luv(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
+    {\r
+        #if (CUDA_VERSION < 5000)\r
+            (void)src;\r
+            (void)dst;\r
+            (void)dcn;\r
+            (void)stream;\r
+            CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );\r
+        #else\r
+            CV_Assert(src.depth() == CV_8U);\r
+            CV_Assert(src.channels() == 3 || src.channels() == 4);\r
+\r
+            dcn = src.channels();\r
+\r
+            dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn));\r
+\r
+            NppStreamHandler h(StreamAccessor::getStream(stream));\r
+\r
+            NppiSize oSizeROI;\r
+            oSizeROI.width = src.cols;\r
+            oSizeROI.height = src.rows;\r
+\r
+            if (dcn == 3)\r
+                nppSafeCall( nppiRGBToLUV_8u_C3R(src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI) );\r
+            else\r
+                nppSafeCall( nppiRGBToLUV_8u_AC4R(src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI) );\r
+        #endif\r
+    }\r
+\r
+    void luv_to_rgb(const GpuMat& src, GpuMat& dst, int dcn, Stream& stream)\r
+    {\r
+        #if (CUDA_VERSION < 5000)\r
+            (void)src;\r
+            (void)dst;\r
+            (void)dcn;\r
+            (void)stream;\r
+            CV_Error( CV_StsBadFlag, "Unknown/unsupported color conversion code" );\r
+        #else\r
+            CV_Assert(src.depth() == CV_8U);\r
+            CV_Assert(src.channels() == 3 || src.channels() == 4);\r
+\r
+            dcn = src.channels();\r
+\r
+            dst.create(src.size(), CV_MAKETYPE(src.depth(), dcn));\r
+\r
+            NppStreamHandler h(StreamAccessor::getStream(stream));\r
+\r
+            NppiSize oSizeROI;\r
+            oSizeROI.width = src.cols;\r
+            oSizeROI.height = src.rows;\r
+\r
+            if (dcn == 3)\r
+                nppSafeCall( nppiLUVToRGB_8u_C3R(src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI) );\r
+            else\r
+                nppSafeCall( nppiLUVToRGB_8u_AC4R(src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI) );\r
+        #endif\r
+    }\r
 }\r
 \r
 void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream& stream)\r
@@ -1203,7 +1314,7 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream
         0,                      //                =42\r
         0,                      //                =43\r
 \r
-        0,                      // CV_BGR2Lab     =44\r
+        bgr_to_lab,             // CV_BGR2Lab     =44\r
         0,                      // CV_RGB2Lab     =45\r
 \r
         0,                      // CV_BayerBG2BGR =46\r
@@ -1212,7 +1323,7 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream
         0,                      // CV_BayerGR2BGR =49\r
 \r
         0,                      // CV_BGR2Luv     =50\r
-        0,                      // CV_RGB2Luv     =51\r
+        rgb_to_luv,             // CV_RGB2Luv     =51\r
 \r
         bgr_to_hls,             // CV_BGR2HLS     =52\r
         rgb_to_hls,             // CV_RGB2HLS     =53\r
@@ -1220,10 +1331,10 @@ void cv::gpu::cvtColor(const GpuMat& src, GpuMat& dst, int code, int dcn, Stream
         hsv_to_bgr,             // CV_HSV2BGR     =54\r
         hsv_to_rgb,             // CV_HSV2RGB     =55\r
 \r
-        0,                      // CV_Lab2BGR     =56\r
+        lab_to_bgr,             // CV_Lab2BGR     =56\r
         0,                      // CV_Lab2RGB     =57\r
         0,                      // CV_Luv2BGR     =58\r
-        0,                      // CV_Luv2RGB     =59\r
+        luv_to_rgb,             // CV_Luv2RGB     =59\r
 \r
         hls_to_bgr,             // CV_HLS2BGR     =60\r
         hls_to_rgb,             // CV_HLS2RGB     =61\r
@@ -1292,4 +1403,45 @@ void cv::gpu::swapChannels(GpuMat& image, const int dstOrder[4], Stream& s)
         cudaSafeCall( cudaDeviceSynchronize() );\r
 }\r
 \r
+void cv::gpu::gammaCorrection(const GpuMat& src, GpuMat& dst, bool forward, Stream& stream)\r
+{\r
+#if (CUDA_VERSION < 5000)\r
+    (void)src;\r
+    (void)dst;\r
+    (void)forward;\r
+    (void)stream;\r
+    CV_Error( CV_StsNotImplemented, "This function works only with CUDA 5.0 or higher" );\r
+#else\r
+    typedef NppStatus (*func_t)(const Npp8u* pSrc, int nSrcStep, Npp8u* pDst, int nDstStep, NppiSize oSizeROI);\r
+    typedef NppStatus (*func_inplace_t)(Npp8u* pSrcDst, int nSrcDstStep, NppiSize oSizeROI);\r
+\r
+    static const func_t funcs[2][5] =\r
+    {\r
+        {0, 0, 0, nppiGammaInv_8u_C3R, nppiGammaInv_8u_AC4R},\r
+        {0, 0, 0, nppiGammaFwd_8u_C3R, nppiGammaFwd_8u_AC4R}\r
+    };\r
+    static const func_inplace_t funcs_inplace[2][5] =\r
+    {\r
+        {0, 0, 0, nppiGammaInv_8u_C3IR, nppiGammaInv_8u_AC4IR},\r
+        {0, 0, 0, nppiGammaFwd_8u_C3IR, nppiGammaFwd_8u_AC4IR}\r
+    };\r
+\r
+    CV_Assert(src.type() == CV_8UC3 || src.type() == CV_8UC4);\r
+\r
+    dst.create(src.size(), src.type());\r
+\r
+    NppStreamHandler h(StreamAccessor::getStream(stream));\r
+\r
+    NppiSize oSizeROI;\r
+    oSizeROI.width = src.cols;\r
+    oSizeROI.height = src.rows;\r
+\r
+    if (dst.data == src.data)\r
+        funcs_inplace[forward][src.channels()](dst.ptr<Npp8u>(), static_cast<int>(src.step), oSizeROI);\r
+    else\r
+        funcs[forward][src.channels()](src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI);\r
+\r
+#endif\r
+}\r
+\r
 #endif /* !defined (HAVE_CUDA) */\r

diff --git a/modules/gpu/test/test_color.cpp b/modules/gpu/test/test_color.cpp
index 1d3ced0..1c08d6a 100644 (file)
--- 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,