added min/max filters to gpu module.

added supports of 4-channels image to gpu::minMax and gpu::cvtColor for RGB <-> YCrCB, RGB <-> YUV and RGB <-> XYZ color conversion.

diff --git a/modules/gpu/include/opencv2/gpu/gpu.hpp b/modules/gpu/include/opencv2/gpu/gpu.hpp
index e2bb5ded3c967d47d78b294b8c797d7f29337d29..7d0074d5fd2573c90e605af242fec244f2550747 100644 (file)
--- a/modules/gpu/include/opencv2/gpu/gpu.hpp
+++ b/modules/gpu/include/opencv2/gpu/gpu.hpp
@@ -411,7 +411,7 @@ namespace cv
         CV_EXPORTS Scalar sum(const GpuMat& m);\r
 \r
         //! finds global minimum and maximum array elements and returns their values\r
-        //! supports only CV_8UC1 type\r
+        //! supports CV_8UC1 and CV_8UC4 type\r
         //! disabled until fix npp bug\r
         CV_EXPORTS void minMax(const GpuMat& src, double* minVal, double* maxVal = 0);\r
 \r
@@ -649,6 +649,12 @@ namespace cv
         //! returns the Gaussian filter engine\r
         CV_EXPORTS Ptr<FilterEngine_GPU> createGaussianFilter_GPU(int type, Size ksize, double sigma1, double sigma2 = 0);\r
 \r
+        //! returns maximum filter\r
+        CV_EXPORTS Ptr<BaseFilter_GPU> getMaxFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor = Point(-1,-1));\r
+\r
+        //! returns minimum filter\r
+        CV_EXPORTS Ptr<BaseFilter_GPU> getMinFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor = Point(-1,-1));\r
+\r
         //! smooths the image using the normalized box filter\r
         //! supports CV_8UC1, CV_8UC4 types\r
         CV_EXPORTS void boxFilter(const GpuMat& src, GpuMat& dst, int ddepth, Size ksize, Point anchor = Point(-1,-1));\r

diff --git a/modules/gpu/src/arithm.cpp b/modules/gpu/src/arithm.cpp
index 4104cb0ab1567081142c838efbd24ea8095a3b18..e5eef84fb8794f7e7b18d21f2de679c6a081dd46 100644 (file)
--- a/modules/gpu/src/arithm.cpp
+++ b/modules/gpu/src/arithm.cpp
@@ -419,24 +419,55 @@ Scalar cv::gpu::sum(const GpuMat& src)
 ////////////////////////////////////////////////////////////////////////\r
 // minMax\r
 \r
-void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal) \r
+namespace\r
 {\r
-    CV_Assert(!"disabled until fix npp bug");\r
-    CV_Assert(src.type() == CV_8UC1);\r
+    void minMax_c1(const GpuMat& src, double* minVal, double* maxVal)\r
+    {\r
+        NppiSize sz;\r
+        sz.width  = src.cols;\r
+        sz.height = src.rows;\r
 \r
-    NppiSize sz;\r
-    sz.width  = src.cols;\r
-    sz.height = src.rows;\r
+        Npp8u min_res, max_res;\r
+\r
+        nppSafeCall( nppiMinMax_8u_C1R(src.ptr<Npp8u>(), src.step, sz, &min_res, &max_res) );\r
 \r
-    Npp8u min_res, max_res;\r
+        if (minVal)\r
+            *minVal = min_res;\r
 \r
-    nppSafeCall( nppiMinMax_8u_C1R(src.ptr<Npp8u>(), src.step, sz, &min_res, &max_res) );\r
+        if (maxVal)\r
+            *maxVal = max_res;\r
+    }\r
+\r
+    void minMax_c4(const GpuMat& src, double* minVal, double* maxVal)\r
+    {\r
+        NppiSize sz;\r
+        sz.width  = src.cols;\r
+        sz.height = src.rows;\r
 \r
-    if (minVal)\r
-        *minVal = min_res;\r
+        Npp8u* cuMin = nppsMalloc_8u(4);\r
+        Npp8u* cuMax = nppsMalloc_8u(4);\r
+\r
+        nppSafeCall( nppiMinMax_8u_C4R(src.ptr<Npp8u>(), src.step, sz, cuMin, cuMax) );\r
+\r
+        if (minVal)\r
+            cudaMemcpy(minVal, cuMin, 4 * sizeof(Npp8u), cudaMemcpyDeviceToHost);        \r
+        if (maxVal)\r
+            cudaMemcpy(maxVal, cuMax, 4 * sizeof(Npp8u), cudaMemcpyDeviceToHost);\r
+\r
+        nppsFree(cuMin);\r
+        nppsFree(cuMax);\r
+    }\r
+}\r
+\r
+void cv::gpu::minMax(const GpuMat& src, double* minVal, double* maxVal) \r
+{\r
+    typedef void (*minMax_t)(const GpuMat& src, double* minVal, double* maxVal);\r
+    static const minMax_t minMax_callers[] = {0, minMax_c1, 0, 0, minMax_c4};\r
+\r
+    CV_Assert(!"disabled until fix npp bug");\r
+    CV_Assert(src.type() == CV_8UC1 || src.type() == CV_8UC4);\r
 \r
-    if (maxVal)\r
-        *maxVal = max_res;\r
+    minMax_callers[src.channels()](src, minVal, maxVal);\r
 }\r
 \r
 ////////////////////////////////////////////////////////////////////////\r

diff --git a/modules/gpu/src/cuda/color.cu b/modules/gpu/src/cuda/color.cu
index 3b81eb6a941f30921e5af9db7ca2fcad40ebf920..0384b1dbf5c74c03bdac907ed526513f5694b5ad 100644 (file)
--- a/modules/gpu/src/cuda/color.cu
+++ b/modules/gpu/src/cuda/color.cu
@@ -625,9 +625,8 @@ namespace imgproc
 \r
     template <typename T> struct RGB2YCrCbConverter \r
     {\r
-        typedef typename TypeVec<T, 3>::vec_t dst_t;\r
-\r
-        static __device__ void cvt(const T* src, dst_t& dst, int bidx)\r
+        template <typename D>\r
+        static __device__ void cvt(const T* src, D& dst, int bidx)\r
         {\r
             const int delta = ColorChannel<T>::half() * (1 << yuv_shift);\r
 \r
@@ -642,9 +641,8 @@ namespace imgproc
     };\r
     template<> struct RGB2YCrCbConverter<float>\r
     {\r
-        typedef typename TypeVec<float, 3>::vec_t dst_t;\r
-\r
-        static __device__ void cvt(const float* src, dst_t& dst, int bidx)\r
+        template <typename D>\r
+        static __device__ void cvt(const float* src, D& dst, int bidx)\r
         {\r
             dst.x = src[0] * cYCrCbCoeffs_f[0] + src[1] * cYCrCbCoeffs_f[1] + src[2] * cYCrCbCoeffs_f[2];\r
             dst.y = (src[bidx^2] - dst.x) * cYCrCbCoeffs_f[3] + ColorChannel<float>::half();\r
@@ -652,11 +650,11 @@ namespace imgproc
         }\r
     };\r
 \r
-    template <int SRCCN, typename T>\r
+    template <int SRCCN, int DSTCN, typename T>\r
     __global__ void RGB2YCrCb(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx)\r
     {\r
         typedef typename TypeVec<T, SRCCN>::vec_t src_t;\r
-        typedef typename TypeVec<T, 3>::vec_t dst_t;\r
+        typedef typename TypeVec<T, DSTCN>::vec_t dst_t;\r
 \r
                const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
                const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
@@ -668,30 +666,28 @@ namespace imgproc
 \r
             RGB2YCrCbConverter<T>::cvt(((const T*)(&src)), dst, bidx);\r
             \r
-            *(dst_t*)(dst_ + y * dst_step + x * 3 * sizeof(T)) = dst;\r
+            *(dst_t*)(dst_ + y * dst_step + x * DSTCN * sizeof(T)) = dst;\r
         }\r
     }\r
 \r
-    template <typename T> struct YCrCb2RGBConvertor\r
+    template <typename D> struct YCrCb2RGBConvertor\r
     {\r
-        typedef typename TypeVec<T, 3>::vec_t src_t;\r
-\r
-        static __device__ void cvt(const src_t& src, T* dst, int bidx)\r
+        template <typename T>\r
+        static __device__ void cvt(const T& src, D* dst, int bidx)\r
         {\r
-            const int b = src.x + CV_DESCALE((src.z - ColorChannel<T>::half()) * cYCrCbCoeffs_i[3], yuv_shift);\r
-            const int g = src.x + CV_DESCALE((src.z - ColorChannel<T>::half()) * cYCrCbCoeffs_i[2] + (src.y - ColorChannel<T>::half()) * cYCrCbCoeffs_i[1], yuv_shift);\r
-            const int r = src.x + CV_DESCALE((src.y - ColorChannel<T>::half()) * cYCrCbCoeffs_i[0], yuv_shift);\r
+            const int b = src.x + CV_DESCALE((src.z - ColorChannel<D>::half()) * cYCrCbCoeffs_i[3], yuv_shift);\r
+            const int g = src.x + CV_DESCALE((src.z - ColorChannel<D>::half()) * cYCrCbCoeffs_i[2] + (src.y - ColorChannel<D>::half()) * cYCrCbCoeffs_i[1], yuv_shift);\r
+            const int r = src.x + CV_DESCALE((src.y - ColorChannel<D>::half()) * cYCrCbCoeffs_i[0], yuv_shift);\r
 \r
-            dst[bidx] = saturate_cast<T>(b);\r
-            dst[1] = saturate_cast<T>(g);\r
-            dst[bidx^2] = saturate_cast<T>(r);\r
+            dst[bidx] = saturate_cast<D>(b);\r
+            dst[1] = saturate_cast<D>(g);\r
+            dst[bidx^2] = saturate_cast<D>(r);\r
         }\r
     };\r
     template <> struct YCrCb2RGBConvertor<float>\r
     {\r
-        typedef typename TypeVec<float, 3>::vec_t src_t;\r
-\r
-        static __device__ void cvt(const src_t& src, float* dst, int bidx)\r
+        template <typename T>\r
+        static __device__ void cvt(const T& src, float* dst, int bidx)\r
         {\r
             dst[bidx] = src.x + (src.z - ColorChannel<float>::half()) * cYCrCbCoeffs_f[3];\r
             dst[1] = src.x + (src.z - ColorChannel<float>::half()) * cYCrCbCoeffs_f[2] + (src.y - ColorChannel<float>::half()) * cYCrCbCoeffs_f[1];\r
@@ -699,10 +695,10 @@ namespace imgproc
         }\r
     };\r
 \r
-    template <int DSTCN, typename T>\r
+    template <int SRCCN, int DSTCN, typename T>\r
     __global__ void YCrCb2RGB(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols, int bidx)\r
     {\r
-        typedef typename TypeVec<T, 3>::vec_t src_t;\r
+        typedef typename TypeVec<T, SRCCN>::vec_t src_t;\r
         typedef typename TypeVec<T, DSTCN>::vec_t dst_t;\r
 \r
                const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
@@ -710,7 +706,7 @@ namespace imgproc
 \r
         if (y < rows && x < cols)\r
         {\r
-            src_t src = *(const src_t*)(src_ + y * src_step + x * 3 * sizeof(T));\r
+            src_t src = *(const src_t*)(src_ + y * src_step + x * SRCCN * sizeof(T));\r
             dst_t dst;\r
 \r
             YCrCb2RGBConvertor<T>::cvt(src, ((T*)(&dst)), bidx);\r
@@ -723,7 +719,7 @@ namespace imgproc
 \r
 namespace cv { namespace gpu { namespace improc\r
 {\r
-    template <typename T, int SRCCN>\r
+    template <typename T, int SRCCN, int DSTCN>\r
     void RGB2YCrCb_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)\r
     {\r
         dim3 threads(32, 8, 1);\r
@@ -732,53 +728,56 @@ namespace cv { namespace gpu { namespace improc
         grid.x = divUp(src.cols, threads.x);\r
         grid.y = divUp(src.rows, threads.y);\r
 \r
-        imgproc::RGB2YCrCb<SRCCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+        imgproc::RGB2YCrCb<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
             dst.ptr, dst.step, src.rows, src.cols, bidx);\r
 \r
         if (stream == 0)\r
             cudaSafeCall( cudaThreadSynchronize() );\r
     }\r
 \r
-    void RGB2YCrCb_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, const int* coeffs, cudaStream_t stream)\r
+    void RGB2YCrCb_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*RGB2YCrCb_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
-        static const RGB2YCrCb_caller_t RGB2YCrCb_callers[2] = \r
+        static const RGB2YCrCb_caller_t RGB2YCrCb_callers[2][2] = \r
         {\r
-            RGB2YCrCb_caller<uchar, 3>, RGB2YCrCb_caller<uchar, 4>\r
+            {RGB2YCrCb_caller<uchar, 3, 3>, RGB2YCrCb_caller<uchar, 3, 4>},\r
+            {RGB2YCrCb_caller<uchar, 4, 3>, RGB2YCrCb_caller<uchar, 4, 4>}\r
         };\r
 \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) );\r
 \r
-        RGB2YCrCb_callers[srccn-3](src, dst, bidx, stream);\r
+        RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
     }\r
 \r
-    void RGB2YCrCb_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, const int* coeffs, cudaStream_t stream)\r
+    void RGB2YCrCb_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*RGB2YCrCb_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
-        static const RGB2YCrCb_caller_t RGB2YCrCb_callers[2] = \r
+        static const RGB2YCrCb_caller_t RGB2YCrCb_callers[2][2] = \r
         {\r
-            RGB2YCrCb_caller<unsigned short, 3>, RGB2YCrCb_caller<unsigned short, 4>\r
+            {RGB2YCrCb_caller<unsigned short, 3, 3>, RGB2YCrCb_caller<unsigned short, 3, 4>},\r
+            {RGB2YCrCb_caller<unsigned short, 4, 3>, RGB2YCrCb_caller<unsigned short, 4, 4>}\r
         };\r
         \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cYCrCbCoeffs_i, coeffs, 5 * sizeof(int)) );\r
 \r
-        RGB2YCrCb_callers[srccn-3](src, dst, bidx, stream);\r
+        RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
     }\r
 \r
-    void RGB2YCrCb_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, const float* coeffs, cudaStream_t stream)\r
+    void RGB2YCrCb_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const float* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*RGB2YCrCb_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
-        static const RGB2YCrCb_caller_t RGB2YCrCb_callers[2] = \r
+        static const RGB2YCrCb_caller_t RGB2YCrCb_callers[2][2] = \r
         {\r
-            RGB2YCrCb_caller<float, 3>, RGB2YCrCb_caller<float, 4>\r
+            {RGB2YCrCb_caller<float, 3, 3>, RGB2YCrCb_caller<float, 3, 4>},\r
+            {RGB2YCrCb_caller<float, 4, 3>, RGB2YCrCb_caller<float, 4, 4>}\r
         };\r
         \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cYCrCbCoeffs_f, coeffs, 5 * sizeof(float)) );\r
 \r
-        RGB2YCrCb_callers[srccn-3](src, dst, bidx, stream);\r
+        RGB2YCrCb_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
     }\r
     \r
-    template <typename T, int DSTCN>\r
+    template <typename T, int SRCCN, int DSTCN>\r
     void YCrCb2RGB_caller(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream)\r
     {\r
         dim3 threads(32, 8, 1);\r
@@ -787,50 +786,53 @@ namespace cv { namespace gpu { namespace improc
         grid.x = divUp(src.cols, threads.x);\r
         grid.y = divUp(src.rows, threads.y);\r
 \r
-        imgproc::YCrCb2RGB<DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+        imgproc::YCrCb2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
             dst.ptr, dst.step, src.rows, src.cols, bidx);\r
 \r
         if (stream == 0)\r
             cudaSafeCall( cudaThreadSynchronize() );\r
     }\r
 \r
-    void YCrCb2RGB_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream)\r
+    void YCrCb2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*YCrCb2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
-        static const YCrCb2RGB_caller_t YCrCb2RGB_callers[2] = \r
+        static const YCrCb2RGB_caller_t YCrCb2RGB_callers[2][2] = \r
         {\r
-            YCrCb2RGB_caller<uchar, 3>, YCrCb2RGB_caller<uchar, 4>\r
+            {YCrCb2RGB_caller<uchar, 3, 3>, YCrCb2RGB_caller<uchar, 3, 4>},\r
+            {YCrCb2RGB_caller<uchar, 4, 3>, YCrCb2RGB_caller<uchar, 4, 4>}\r
         };\r
 \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) );\r
 \r
-        YCrCb2RGB_callers[dstcn-3](src, dst, bidx, stream);\r
+        YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
     }\r
 \r
-    void YCrCb2RGB_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream)\r
+    void YCrCb2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*YCrCb2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
-        static const YCrCb2RGB_caller_t YCrCb2RGB_callers[2] = \r
+        static const YCrCb2RGB_caller_t YCrCb2RGB_callers[2][2] = \r
         {\r
-            YCrCb2RGB_caller<unsigned short, 3>, YCrCb2RGB_caller<unsigned short, 4>\r
+            {YCrCb2RGB_caller<unsigned short, 3, 3>, YCrCb2RGB_caller<unsigned short, 3, 4>},\r
+            {YCrCb2RGB_caller<unsigned short, 4, 3>, YCrCb2RGB_caller<unsigned short, 4, 4>}\r
         };\r
         \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cYCrCbCoeffs_i, coeffs, 4 * sizeof(int)) );\r
 \r
-        YCrCb2RGB_callers[dstcn-3](src, dst, bidx, stream);\r
+        YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
     }\r
 \r
-    void YCrCb2RGB_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int dstcn, int bidx, const float* coeffs, cudaStream_t stream)\r
+    void YCrCb2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const float* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*YCrCb2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
-        static const YCrCb2RGB_caller_t YCrCb2RGB_callers[2] = \r
+        static const YCrCb2RGB_caller_t YCrCb2RGB_callers[2][2] = \r
         {\r
-            YCrCb2RGB_caller<float, 3>, YCrCb2RGB_caller<float, 4>\r
+            {YCrCb2RGB_caller<float, 3, 3>, YCrCb2RGB_caller<float, 3, 4>},\r
+            {YCrCb2RGB_caller<float, 4, 3>, YCrCb2RGB_caller<float, 4, 4>}\r
         };\r
         \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cYCrCbCoeffs_f, coeffs, 4 * sizeof(float)) );\r
 \r
-        YCrCb2RGB_callers[dstcn-3](src, dst, bidx, stream);\r
+        YCrCb2RGB_callers[srccn-3][dstcn-3](src, dst, bidx, stream);\r
     }\r
 }}}\r
 \r
@@ -843,38 +845,30 @@ namespace imgproc
 \r
     template <typename T> struct RGB2XYZConvertor\r
     {\r
-        typedef typename TypeVec<T, 3>::vec_t dst_t;\r
-        static __device__ dst_t cvt(const T* src)\r
+        template <typename D>\r
+        static __device__ void cvt(const T* src, D& dst)\r
         {\r
-            dst_t dst;\r
-\r
                dst.x = saturate_cast<T>(CV_DESCALE(src[0] * cXYZ_D65i[0] + src[1] * cXYZ_D65i[1] + src[2] * cXYZ_D65i[2], xyz_shift));\r
                dst.y = saturate_cast<T>(CV_DESCALE(src[0] * cXYZ_D65i[3] + src[1] * cXYZ_D65i[4] + src[2] * cXYZ_D65i[5], xyz_shift));\r
                dst.z = saturate_cast<T>(CV_DESCALE(src[0] * cXYZ_D65i[6] + src[1] * cXYZ_D65i[7] + src[2] * cXYZ_D65i[8], xyz_shift));\r
-\r
-            return dst;\r
         }\r
     };\r
     template <> struct RGB2XYZConvertor<float>\r
     {\r
-        typedef typename TypeVec<float, 3>::vec_t dst_t;\r
-        static __device__ dst_t cvt(const float* src)\r
+        template <typename D>\r
+        static __device__ void cvt(const float* src, D& dst)\r
         {\r
-            dst_t dst;\r
-\r
                dst.x = src[0] * cXYZ_D65f[0] + src[1] * cXYZ_D65f[1] + src[2] * cXYZ_D65f[2];\r
                dst.y = src[0] * cXYZ_D65f[3] + src[1] * cXYZ_D65f[4] + src[2] * cXYZ_D65f[5];\r
                dst.z = src[0] * cXYZ_D65f[6] + src[1] * cXYZ_D65f[7] + src[2] * cXYZ_D65f[8];\r
-\r
-            return dst;\r
         }\r
     };\r
 \r
-    template <int SRCCN, typename T>\r
+    template <int SRCCN, int DSTCN, typename T>\r
     __global__ void RGB2XYZ(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols)\r
     {\r
         typedef typename TypeVec<T, SRCCN>::vec_t src_t;\r
-        typedef typename TypeVec<T, 3>::vec_t dst_t;\r
+        typedef typename TypeVec<T, DSTCN>::vec_t dst_t;\r
 \r
                const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
                const int y = blockDim.y * blockIdx.y + threadIdx.y;\r
@@ -882,25 +876,28 @@ namespace imgproc
         if (y < rows && x < cols)\r
         {\r
             src_t src = *(const src_t*)(src_ + y * src_step + x * SRCCN * sizeof(T));\r
+\r
+            dst_t dst;\r
+            RGB2XYZConvertor<T>::cvt((const T*)(&src), dst);\r
             \r
-            *(dst_t*)(dst_ + y * dst_step + x * 3 * sizeof(T)) = RGB2XYZConvertor<T>::cvt((const T*)(&src));\r
+            *(dst_t*)(dst_ + y * dst_step + x * DSTCN * sizeof(T)) = dst;\r
         }\r
     }\r
 \r
-    template <typename T> struct XYZ2RGBConvertor\r
+    template <typename D> struct XYZ2RGBConvertor\r
     {\r
-        typedef typename TypeVec<T, 3>::vec_t src_t;\r
-        static __device__ void cvt(const src_t& src, T* dst)\r
+        template <typename T>\r
+        static __device__ void cvt(const T& src, D* dst)\r
         {\r
-            dst[0] = saturate_cast<T>(CV_DESCALE(src.x * cXYZ_D65i[0] + src.y * cXYZ_D65i[1] + src.z * cXYZ_D65i[2], xyz_shift));\r
-                   dst[1] = saturate_cast<T>(CV_DESCALE(src.x * cXYZ_D65i[3] + src.y * cXYZ_D65i[4] + src.z * cXYZ_D65i[5], xyz_shift));\r
-                   dst[2] = saturate_cast<T>(CV_DESCALE(src.x * cXYZ_D65i[6] + src.y * cXYZ_D65i[7] + src.z * cXYZ_D65i[8], xyz_shift));\r
+            dst[0] = saturate_cast<D>(CV_DESCALE(src.x * cXYZ_D65i[0] + src.y * cXYZ_D65i[1] + src.z * cXYZ_D65i[2], xyz_shift));\r
+                   dst[1] = saturate_cast<D>(CV_DESCALE(src.x * cXYZ_D65i[3] + src.y * cXYZ_D65i[4] + src.z * cXYZ_D65i[5], xyz_shift));\r
+                   dst[2] = saturate_cast<D>(CV_DESCALE(src.x * cXYZ_D65i[6] + src.y * cXYZ_D65i[7] + src.z * cXYZ_D65i[8], xyz_shift));\r
         }\r
     };\r
     template <> struct XYZ2RGBConvertor<float>\r
     {\r
-        typedef typename TypeVec<float, 3>::vec_t src_t;\r
-        static __device__ void cvt(const src_t& src, float* dst)\r
+        template <typename T>\r
+        static __device__ void cvt(const T& src, float* dst)\r
         {\r
             dst[0] = src.x * cXYZ_D65f[0] + src.y * cXYZ_D65f[1] + src.z * cXYZ_D65f[2];\r
                    dst[1] = src.x * cXYZ_D65f[3] + src.y * cXYZ_D65f[4] + src.z * cXYZ_D65f[5];\r
@@ -908,10 +905,10 @@ namespace imgproc
         }\r
     };\r
 \r
-    template <int DSTCN, typename T>\r
+    template <int SRCCN, int DSTCN, typename T>\r
     __global__ void XYZ2RGB(const uchar* src_, size_t src_step, uchar* dst_, size_t dst_step, int rows, int cols)\r
     {\r
-        typedef typename TypeVec<T, 3>::vec_t src_t;\r
+        typedef typename TypeVec<T, SRCCN>::vec_t src_t;\r
         typedef typename TypeVec<T, DSTCN>::vec_t dst_t;\r
 \r
                const int x = blockDim.x * blockIdx.x + threadIdx.x;\r
@@ -919,7 +916,7 @@ namespace imgproc
 \r
         if (y < rows && x < cols)\r
         {\r
-            src_t src = *(const src_t*)(src_ + y * src_step + x * 3 * sizeof(T));\r
+            src_t src = *(const src_t*)(src_ + y * src_step + x * SRCCN * sizeof(T));\r
 \r
             dst_t dst;\r
             XYZ2RGBConvertor<T>::cvt(src, (T*)(&dst));\r
@@ -932,7 +929,7 @@ namespace imgproc
 \r
 namespace cv { namespace gpu { namespace improc\r
 {\r
-    template <typename T, int SRCCN>\r
+    template <typename T, int SRCCN, int DSTCN>\r
     void RGB2XYZ_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream)\r
     {\r
         dim3 threads(32, 8, 1);\r
@@ -941,44 +938,56 @@ namespace cv { namespace gpu { namespace improc
         grid.x = divUp(src.cols, threads.x);\r
         grid.y = divUp(src.rows, threads.y);\r
 \r
-        imgproc::RGB2XYZ<SRCCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+        imgproc::RGB2XYZ<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
             dst.ptr, dst.step, src.rows, src.cols);\r
 \r
         if (stream == 0)\r
             cudaSafeCall( cudaThreadSynchronize() );\r
     }\r
 \r
-    void RGB2XYZ_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, const int* coeffs, cudaStream_t stream)\r
+    void RGB2XYZ_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*RGB2XYZ_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream);\r
-        static const RGB2XYZ_caller_t RGB2XYZ_callers[] = {RGB2XYZ_caller<uchar, 3>, RGB2XYZ_caller<uchar, 4>};\r
+        static const RGB2XYZ_caller_t RGB2XYZ_callers[2][2] = \r
+        {\r
+            {RGB2XYZ_caller<uchar, 3, 3>, RGB2XYZ_caller<uchar, 3, 4>},\r
+            {RGB2XYZ_caller<uchar, 4, 3>, RGB2XYZ_caller<uchar, 4, 4>}\r
+        };\r
 \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
 \r
-        RGB2XYZ_callers[srccn-3](src, dst, stream);\r
+        RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, stream);\r
     }\r
 \r
-    void RGB2XYZ_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, const int* coeffs, cudaStream_t stream)\r
+    void RGB2XYZ_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*RGB2XYZ_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream);\r
-        static const RGB2XYZ_caller_t RGB2XYZ_callers[] = {RGB2XYZ_caller<unsigned short, 3>, RGB2XYZ_caller<unsigned short, 4>};\r
+        static const RGB2XYZ_caller_t RGB2XYZ_callers[2][2] = \r
+        {\r
+            {RGB2XYZ_caller<unsigned short, 3, 3>, RGB2XYZ_caller<unsigned short, 3, 4>},\r
+            {RGB2XYZ_caller<unsigned short, 4, 3>, RGB2XYZ_caller<unsigned short, 4, 4>}\r
+        };\r
         \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
 \r
-        RGB2XYZ_callers[srccn-3](src, dst, stream);\r
+        RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, stream);\r
     }\r
 \r
-    void RGB2XYZ_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, const float* coeffs, cudaStream_t stream)\r
+    void RGB2XYZ_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const float* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*RGB2XYZ_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream);\r
-        static const RGB2XYZ_caller_t RGB2XYZ_callers[] = {RGB2XYZ_caller<float, 3>, RGB2XYZ_caller<float, 4>};\r
+        static const RGB2XYZ_caller_t RGB2XYZ_callers[2][2] = \r
+        {\r
+            {RGB2XYZ_caller<float, 3, 3>, RGB2XYZ_caller<float, 3, 4>},\r
+            {RGB2XYZ_caller<float, 4, 3>, RGB2XYZ_caller<float, 4, 4>}\r
+        };\r
         \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cXYZ_D65f, coeffs, 9 * sizeof(float)) );\r
 \r
-        RGB2XYZ_callers[srccn-3](src, dst, stream);\r
+        RGB2XYZ_callers[srccn-3][dstcn-3](src, dst, stream);\r
     }\r
     \r
-    template <typename T, int DSTCN>\r
+    template <typename T, int SRCCN, int DSTCN>\r
     void XYZ2RGB_caller(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream)\r
     {\r
         dim3 threads(32, 8, 1);\r
@@ -987,41 +996,53 @@ namespace cv { namespace gpu { namespace improc
         grid.x = divUp(src.cols, threads.x);\r
         grid.y = divUp(src.rows, threads.y);\r
 \r
-        imgproc::XYZ2RGB<DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
+        imgproc::XYZ2RGB<SRCCN, DSTCN, T><<<grid, threads, 0, stream>>>(src.ptr, src.step, \r
             dst.ptr, dst.step, src.rows, src.cols);\r
 \r
         if (stream == 0)\r
             cudaSafeCall( cudaThreadSynchronize() );\r
     }\r
 \r
-    void XYZ2RGB_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream)\r
+    void XYZ2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*XYZ2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream);\r
-        static const XYZ2RGB_caller_t XYZ2RGB_callers[] = {XYZ2RGB_caller<uchar, 3>, XYZ2RGB_caller<uchar, 4>};\r
+        static const XYZ2RGB_caller_t XYZ2RGB_callers[2][2] = \r
+        {\r
+            {XYZ2RGB_caller<uchar, 3, 3>, XYZ2RGB_caller<uchar, 3, 4>},\r
+            {XYZ2RGB_caller<uchar, 4, 3>, XYZ2RGB_caller<uchar, 4, 4>}\r
+        };\r
 \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
 \r
-        XYZ2RGB_callers[dstcn-3](src, dst, stream);\r
+        XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, stream);\r
     }\r
 \r
-    void XYZ2RGB_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream)\r
+    void XYZ2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*XYZ2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream);\r
-        static const XYZ2RGB_caller_t XYZ2RGB_callers[] = {XYZ2RGB_caller<unsigned short, 3>, XYZ2RGB_caller<unsigned short, 4>};\r
+        static const XYZ2RGB_caller_t XYZ2RGB_callers[2][2] = \r
+        {\r
+            {XYZ2RGB_caller<unsigned short, 3, 3>, XYZ2RGB_caller<unsigned short, 3, 4>},\r
+            {XYZ2RGB_caller<unsigned short, 4, 3>, XYZ2RGB_caller<unsigned short, 4, 4>}\r
+        };\r
         \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cXYZ_D65i, coeffs, 9 * sizeof(int)) );\r
 \r
-        XYZ2RGB_callers[dstcn-3](src, dst, stream);\r
+        XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, stream);\r
     }\r
 \r
-    void XYZ2RGB_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int dstcn, const float* coeffs, cudaStream_t stream)\r
+    void XYZ2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const float* coeffs, cudaStream_t stream)\r
     {\r
         typedef void (*XYZ2RGB_caller_t)(const DevMem2D& src, const DevMem2D& dst, cudaStream_t stream);\r
-        static const XYZ2RGB_caller_t XYZ2RGB_callers[] = {XYZ2RGB_caller<float, 3>, XYZ2RGB_caller<float, 4>};\r
+        static const XYZ2RGB_caller_t XYZ2RGB_callers[2][2] = \r
+        {\r
+            {XYZ2RGB_caller<float, 3, 3>, XYZ2RGB_caller<float, 3, 4>},\r
+            {XYZ2RGB_caller<float, 4, 3>, XYZ2RGB_caller<float, 4, 4>}\r
+        };\r
         \r
         cudaSafeCall( cudaMemcpyToSymbol(imgproc::cXYZ_D65f, coeffs, 9 * sizeof(float)) );\r
 \r
-        XYZ2RGB_callers[dstcn-3](src, dst, stream);\r
+        XYZ2RGB_callers[srccn-3][dstcn-3](src, dst, stream);\r
     }\r
 }}}\r
 \r

diff --git a/modules/gpu/src/filtering_npp.cpp b/modules/gpu/src/filtering_npp.cpp
index 1ca51e152162761b51d1a3131ed6f3e428424606..97f08a7606156968cca401d06e668140b066c3b3 100644 (file)
--- a/modules/gpu/src/filtering_npp.cpp
+++ b/modules/gpu/src/filtering_npp.cpp
@@ -63,6 +63,9 @@ Ptr<BaseColumnFilter_GPU> cv::gpu::getLinearColumnFilter_GPU(int, int, const Gpu
 Ptr<FilterEngine_GPU> cv::gpu::createSeparableLinearFilter_GPU(int, int, const Mat&, const Mat&, const Point&) { throw_nogpu(); return Ptr<FilterEngine_GPU>(0); }\r
 Ptr<FilterEngine_GPU> cv::gpu::createDerivFilter_GPU(int, int, int, int, int) { throw_nogpu(); return Ptr<FilterEngine_GPU>(0); }\r
 Ptr<FilterEngine_GPU> cv::gpu::createGaussianFilter_GPU(int, Size, double, double) { throw_nogpu(); return Ptr<FilterEngine_GPU>(0); }\r
+Ptr<BaseFilter_GPU> cv::gpu::getMaxFilter_GPU(int, int, const Size&, Point) { throw_nogpu(); return Ptr<BaseFilter_GPU>(0); }\r
+Ptr<BaseFilter_GPU> cv::gpu::getMinFilter_GPU(int, int, const Size&, Point) { throw_nogpu(); return Ptr<BaseFilter_GPU>(0); }\r
+\r
 void cv::gpu::boxFilter(const GpuMat&, GpuMat&, int, Size, Point) { throw_nogpu(); }\r
 void cv::gpu::erode( const GpuMat&, GpuMat&, const Mat&, Point, int) { throw_nogpu(); }\r
 void cv::gpu::dilate( const GpuMat&, GpuMat&, const Mat&, Point, int) { throw_nogpu(); }\r
@@ -105,20 +108,20 @@ namespace
         int scale = nDivisor && (kernel.depth() == CV_32F || kernel.depth() == CV_64F) ? 256 : 1;\r
         if (nDivisor) *nDivisor = scale;\r
         \r
-        Mat cont_krnl = (kernel.isContinuous() ? kernel : kernel.clone()).reshape(1, 1);\r
-        Mat temp;\r
-        cont_krnl.convertTo(temp, type, scale);\r
+        Mat temp(kernel.size(), type);\r
+        kernel.convertTo(temp, type, scale);\r
+        Mat cont_krnl = temp.reshape(1, 1);\r
 \r
         if (reverse)\r
         {\r
-            int count = temp.cols >> 1;\r
+            int count = cont_krnl.cols >> 1;\r
             for (int i = 0; i < count; ++i)\r
             {\r
-                std::swap(temp.at<int>(0, i), temp.at<int>(0, temp.cols - 1 - i));\r
+                std::swap(cont_krnl.at<int>(0, i), cont_krnl.at<int>(0, cont_krnl.cols - 1 - i));\r
             }\r
         }\r
 \r
-        gpu_krnl.upload(temp);\r
+        gpu_krnl.upload(cont_krnl);\r
     } \r
 }\r
 \r
@@ -785,4 +788,58 @@ void cv::gpu::GaussianBlur(const GpuMat& src, GpuMat& dst, Size ksize, double si
     f->apply(src, dst);\r
 }\r
 \r
+////////////////////////////////////////////////////////////////////////////////////////////////////\r
+// Image Rank Filter\r
+\r
+namespace\r
+{\r
+    typedef NppStatus (*nppFilterRank_t)(const Npp8u * pSrc, Npp32s nSrcStep, Npp8u * pDst, Npp32s nDstStep, NppiSize oSizeROI, \r
+        NppiSize oMaskSize, NppiPoint oAnchor);\r
+\r
+    class NPPRankFilter : public BaseFilter_GPU\r
+    {\r
+    public:\r
+        NPPRankFilter(const Size& ksize_, const Point& anchor_, nppFilterRank_t func_) : BaseFilter_GPU(ksize_, anchor_), func(func_) {}\r
+\r
+        virtual void operator()(const GpuMat& src, GpuMat& dst)\r
+        {\r
+            NppiSize sz;\r
+            sz.width = src.cols;\r
+            sz.height = src.rows;\r
+            NppiSize oKernelSize;\r
+            oKernelSize.height = ksize.height;\r
+            oKernelSize.width = ksize.width;\r
+            NppiPoint oAnchor;\r
+            oAnchor.x = anchor.x;\r
+            oAnchor.y = anchor.y;\r
+            \r
+            nppSafeCall( func(src.ptr<Npp8u>(), src.step, dst.ptr<Npp8u>(), dst.step, sz, oKernelSize, oAnchor) );\r
+        }\r
+\r
+        nppFilterRank_t func;\r
+    };\r
+}\r
+\r
+Ptr<BaseFilter_GPU> cv::gpu::getMaxFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor)\r
+{\r
+    static const nppFilterRank_t nppFilterRank_callers[] = {0, nppiFilterMax_8u_C1R, 0, 0, nppiFilterMax_8u_C4R};\r
+\r
+    CV_Assert((srcType == CV_8UC1 || srcType == CV_8UC4) && dstType == srcType); \r
+\r
+    normalizeAnchor(anchor, ksize);\r
+\r
+    return Ptr<BaseFilter_GPU>(new NPPRankFilter(ksize, anchor, nppFilterRank_callers[CV_MAT_CN(srcType)]));\r
+}\r
+\r
+Ptr<BaseFilter_GPU> cv::gpu::getMinFilter_GPU(int srcType, int dstType, const Size& ksize, Point anchor)\r
+{\r
+    static const nppFilterRank_t nppFilterRank_callers[] = {0, nppiFilterMin_8u_C1R, 0, 0, nppiFilterMin_8u_C4R};\r
+\r
+    CV_Assert((srcType == CV_8UC1 || srcType == CV_8UC4) && dstType == srcType); \r
+\r
+    normalizeAnchor(anchor, ksize);\r
+\r
+    return Ptr<BaseFilter_GPU>(new NPPRankFilter(ksize, anchor, nppFilterRank_callers[CV_MAT_CN(srcType)]));\r
+}\r
+\r
 #endif\r

diff --git a/modules/gpu/src/imgproc_gpu.cpp b/modules/gpu/src/imgproc_gpu.cpp
index 9c88f3ea7a6d445f6c76846369988a72401d83d4..0badf8b2937123db88ca2af88109dd129adf1f8e 100644 (file)
--- a/modules/gpu/src/imgproc_gpu.cpp
+++ b/modules/gpu/src/imgproc_gpu.cpp
@@ -98,21 +98,21 @@ namespace cv { namespace gpu
         void RGB2Gray_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, cudaStream_t stream);\r
         void RGB5x52Gray_gpu(const DevMem2D& src, int green_bits, const DevMem2D& dst, cudaStream_t stream);\r
 \r
-        void RGB2YCrCb_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, const int* coeffs, cudaStream_t stream);\r
-        void RGB2YCrCb_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, const int* coeffs, cudaStream_t stream);\r
-        void RGB2YCrCb_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int bidx, const float* coeffs, cudaStream_t stream);\r
+        void RGB2YCrCb_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream);\r
+        void RGB2YCrCb_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream);\r
+        void RGB2YCrCb_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const float* coeffs, cudaStream_t stream);\r
 \r
-        void YCrCb2RGB_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream);\r
-        void YCrCb2RGB_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream);\r
-        void YCrCb2RGB_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int dstcn, int bidx, const float* coeffs, cudaStream_t stream);\r
+        void YCrCb2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream);\r
+        void YCrCb2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const int* coeffs, cudaStream_t stream);\r
+        void YCrCb2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, int bidx, const float* coeffs, cudaStream_t stream);\r
 \r
-        void RGB2XYZ_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, const int* coeffs, cudaStream_t stream);\r
-        void RGB2XYZ_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, const int* coeffs, cudaStream_t stream);\r
-        void RGB2XYZ_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, const float* coeffs, cudaStream_t stream);\r
+        void RGB2XYZ_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream);\r
+        void RGB2XYZ_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream);\r
+        void RGB2XYZ_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const float* coeffs, cudaStream_t stream);\r
 \r
-        void XYZ2RGB_gpu_8u(const DevMem2D& src, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream);\r
-        void XYZ2RGB_gpu_16u(const DevMem2D& src, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream);\r
-        void XYZ2RGB_gpu_32f(const DevMem2D& src, const DevMem2D& dst, int dstcn, const float* coeffs, cudaStream_t stream);\r
+        void XYZ2RGB_gpu_8u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream);\r
+        void XYZ2RGB_gpu_16u(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const int* coeffs, cudaStream_t stream);\r
+        void XYZ2RGB_gpu_32f(const DevMem2D& src, int srccn, const DevMem2D& dst, int dstcn, const float* coeffs, cudaStream_t stream);\r
     }\r
 }}\r
 \r
@@ -361,7 +361,8 @@ namespace
             case CV_BGR2YCrCb: case CV_RGB2YCrCb:\r
             case CV_BGR2YUV: case CV_RGB2YUV:\r
                 {\r
-                    CV_Assert( scn == 3 || scn == 4 );\r
+                    if(dcn <= 0) dcn = 3;\r
+                    CV_Assert( (scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) );\r
 \r
                     bidx = code == CV_BGR2YCrCb || code == CV_RGB2YUV ? 0 : 2;\r
 \r
@@ -382,14 +383,14 @@ namespace
                         std::swap(coeffs_i[0], coeffs_i[2]);\r
                     }\r
                         \r
-                    out.create(sz, CV_MAKETYPE(depth, 3));\r
+                    out.create(sz, CV_MAKETYPE(depth, dcn));\r
                     \r
                     if( depth == CV_8U )\r
-                        improc::RGB2YCrCb_gpu_8u(src, scn, out, bidx, coeffs_i, stream);\r
+                        improc::RGB2YCrCb_gpu_8u(src, scn, out, dcn, bidx, coeffs_i, stream);\r
                     else if( depth == CV_16U )\r
-                        improc::RGB2YCrCb_gpu_16u(src, scn, out, bidx, coeffs_i, stream);\r
+                        improc::RGB2YCrCb_gpu_16u(src, scn, out, dcn, bidx, coeffs_i, stream);\r
                     else\r
-                        improc::RGB2YCrCb_gpu_32f(src, scn, out, bidx, coeffs_f, stream);\r
+                        improc::RGB2YCrCb_gpu_32f(src, scn, out, dcn, bidx, coeffs_f, stream);\r
                 }\r
                 break;\r
                 \r
@@ -398,7 +399,7 @@ namespace
                 {\r
                     if (dcn <= 0) dcn = 3;\r
 \r
-                    CV_Assert( scn == 3 && (dcn == 3 || dcn == 4) );\r
+                    CV_Assert( (scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) );\r
 \r
                     bidx = code == CV_YCrCb2BGR || code == CV_YUV2RGB ? 0 : 2;\r
 \r
@@ -414,17 +415,18 @@ namespace
                     out.create(sz, CV_MAKETYPE(depth, dcn));\r
                     \r
                     if( depth == CV_8U )\r
-                        improc::YCrCb2RGB_gpu_8u(src, out, dcn, bidx, coeffs_i, stream);\r
+                        improc::YCrCb2RGB_gpu_8u(src, scn, out, dcn, bidx, coeffs_i, stream);\r
                     else if( depth == CV_16U )\r
-                        improc::YCrCb2RGB_gpu_16u(src, out, dcn, bidx, coeffs_i, stream);\r
+                        improc::YCrCb2RGB_gpu_16u(src, scn, out, dcn, bidx, coeffs_i, stream);\r
                     else\r
-                        improc::YCrCb2RGB_gpu_32f(src, out, dcn, bidx, coeffs_f, stream);\r
+                        improc::YCrCb2RGB_gpu_32f(src, scn, out, dcn, bidx, coeffs_f, stream);\r
                 }\r
                 break;\r
             \r
             case CV_BGR2XYZ: case CV_RGB2XYZ:\r
-                {\r
-                    CV_Assert( scn == 3 || scn == 4 );\r
+                {                    \r
+                    if(dcn <= 0) dcn = 3;\r
+                    CV_Assert( (scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) );\r
 \r
                     bidx = code == CV_BGR2XYZ ? 0 : 2;\r
 \r
@@ -457,21 +459,21 @@ namespace
                         std::swap(coeffs_i[6], coeffs_i[8]);\r
                     }\r
                         \r
-                    out.create(sz, CV_MAKETYPE(depth, 3));\r
+                    out.create(sz, CV_MAKETYPE(depth, dcn));\r
                     \r
                     if( depth == CV_8U )\r
-                        improc::RGB2XYZ_gpu_8u(src, scn, out, coeffs_i, stream);\r
+                        improc::RGB2XYZ_gpu_8u(src, scn, out, dcn, coeffs_i, stream);\r
                     else if( depth == CV_16U )\r
-                        improc::RGB2XYZ_gpu_16u(src, scn, out, coeffs_i, stream);\r
+                        improc::RGB2XYZ_gpu_16u(src, scn, out, dcn, coeffs_i, stream);\r
                     else\r
-                        improc::RGB2XYZ_gpu_32f(src, scn, out, coeffs_f, stream);\r
+                        improc::RGB2XYZ_gpu_32f(src, scn, out, dcn, coeffs_f, stream);\r
                 }\r
                 break;\r
             \r
             case CV_XYZ2BGR: case CV_XYZ2RGB:\r
                 {\r
                     if (dcn <= 0) dcn = 3;\r
-                    CV_Assert( scn == 3 && (dcn == 3 || dcn == 4) );\r
+                    CV_Assert( (scn == 3 || scn == 4) && (dcn == 3 || dcn == 4) );\r
                     bidx = code == CV_XYZ2BGR ? 0 : 2;\r
 \r
                     static const float XYZ2sRGB_D65f[] =\r
@@ -506,11 +508,11 @@ namespace
                     out.create(sz, CV_MAKETYPE(depth, dcn));\r
                     \r
                     if( depth == CV_8U )\r
-                        improc::XYZ2RGB_gpu_8u(src, out, dcn, coeffs_i, stream);\r
+                        improc::XYZ2RGB_gpu_8u(src, scn, out, dcn, coeffs_i, stream);\r
                     else if( depth == CV_16U )\r
-                        improc::XYZ2RGB_gpu_16u(src, out, dcn, coeffs_i, stream);\r
+                        improc::XYZ2RGB_gpu_16u(src, scn, out, dcn, coeffs_i, stream);\r
                     else\r
-                        improc::XYZ2RGB_gpu_32f(src, out, dcn, coeffs_f, stream);\r
+                        improc::XYZ2RGB_gpu_32f(src, scn, out, dcn, coeffs_f, stream);\r
                 }\r
                 break;\r
                 \r