rewrite cuda::cvtColor with new device layer and fix test failures

modules/cudaimgproc/CMakeLists.txt

@@ -6,4 +6,4 @@ set(the_description "CUDA-accelerated Image Processing")
 
 ocv_warnings_disable(CMAKE_CXX_FLAGS /wd4127 /wd4100 /wd4324 /wd4512 /wd4515 -Wundef -Wmissing-declarations -Wshadow -Wunused-parameter)
 
-ocv_define_module(cudaimgproc opencv_imgproc OPTIONAL opencv_cudaarithm opencv_cudafilters)
+ocv_define_module(cudaimgproc opencv_imgproc OPTIONAL opencv_cudev opencv_cudaarithm opencv_cudafilters)

modules/cudev/include/opencv2/cudev/functional/detail/color_cvt.hpp

@@ -156,7 +156,7 @@ namespace color_cvt_detail
             const int g = src.y;
             const int r = bidx == 0 ? src.z : src.x;
             const int a = src.w;
-            return (ushort) ((b >> 3) | ((g & ~7) << 2) | ((r & ~7) << 7) | (a * 0x8000));
+            return (ushort) ((b >> 3) | ((g & ~7) << 2) | ((r & ~7) << 7) | (a ? 0x8000 : 0));
         }
     };
 
@@ -263,7 +263,8 @@ namespace color_cvt_detail
     {
         __device__ ushort operator ()(uchar src) const
         {
-            return (ushort) (src | (src << 5) | (src << 10));
+            const int t = src >> 3;
+            return (ushort)(t | (t << 5) | (t << 10));
         }
     };
 
@@ -272,7 +273,8 @@ namespace color_cvt_detail
     {
         __device__ ushort operator ()(uchar src) const
         {
-            return (ushort) ((src >> 3) | ((src & ~3) << 3) | ((src & ~7) << 8));
+            const int t = src;
+            return (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8));
         }
     };
 
@@ -1227,6 +1229,10 @@ namespace color_cvt_detail
             float G = -0.969256f * X + 1.875991f * Y + 0.041556f * Z;
             float R = 3.240479f * X - 1.537150f * Y - 0.498535f * Z;
 
+            R = ::fminf(::fmaxf(R, 0.f), 1.f);
+            G = ::fminf(::fmaxf(G, 0.f), 1.f);
+            B = ::fminf(::fmaxf(B, 0.f), 1.f);
+
             if (srgb)
             {
                 B = splineInterpolate(B * GAMMA_TAB_SIZE, c_sRGBInvGammaTab, GAMMA_TAB_SIZE);