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

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
 {

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<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(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<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(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<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI) );
+            else
+                nppSafeCall( nppiRGBToLUV_8u_AC4R(src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(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<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI) );
+            else
+                nppSafeCall( nppiLUVToRGB_8u_AC4R(src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(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<Npp8u>(), static_cast<int>(src.step), oSizeROI);
+    else
+        funcs[forward][src.channels()](src.ptr<Npp8u>(), static_cast<int>(src.step), dst.ptr<Npp8u>(), static_cast<int>(dst.step), oSizeROI);
+
+#endif
+}
+
 #endif /* !defined (HAVE_CUDA) */

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,