Added support for YUV2RGB[A]_NV21 and YUV2BGR[A]_NV21 conversion

modules/imgproc/src/color.cpp

@@ -4848,7 +4848,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
     bool ok = false;
     UMat src = _src.getUMat(), dst;
     Size sz = src.size(), dstSz = sz;
-    int scn = src.channels(), depth = src.depth(), bidx;
+    int scn = src.channels(), depth = src.depth(), bidx, uidx;
     int dims = 2, stripeSize = 1;
     ocl::Kernel k;
 
@@ -4960,17 +4960,21 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
                  opts + format("-D dcn=%d -D bidx=%d", dcn, bidx));
         break;
     }
-    case COLOR_YUV2RGB_NV12: case COLOR_YUV2BGR_NV12:
-    case COLOR_YUV2RGBA_NV12: case COLOR_YUV2BGRA_NV12:
+    case COLOR_YUV2RGB_NV12: case COLOR_YUV2BGR_NV12: case COLOR_YUV2RGB_NV21: case COLOR_YUV2BGR_NV21:
+    case COLOR_YUV2RGBA_NV12: case COLOR_YUV2BGRA_NV12: case COLOR_YUV2RGBA_NV21: case COLOR_YUV2BGRA_NV21:
     {
         CV_Assert( scn == 1 );
         CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U );
-        dcn  = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2RGBA_NV12 ? 4 : 3;
-        bidx = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2BGR_NV12 ? 0 : 2;
+        dcn  = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2RGBA_NV12 ||
+               code == COLOR_YUV2BGRA_NV21 || code == COLOR_YUV2RGBA_NV21 ? 4 : 3;
+        bidx = code == COLOR_YUV2BGRA_NV12 || code == COLOR_YUV2BGR_NV12 || 
+               code == COLOR_YUV2BGRA_NV21 || code == COLOR_YUV2BGR_NV21 ? 0 : 2;
+        uidx = code == COLOR_YUV2RGBA_NV21 || code == COLOR_YUV2RGB_NV21 ||
+               code == COLOR_YUV2BGRA_NV21 || code == COLOR_YUV2BGR_NV21 ? 1 : 0;
 
         dstSz = Size(sz.width, sz.height * 2 / 3);
-        k.create("YUV2RGB_NV12", ocl::imgproc::cvtcolor_oclsrc,
-                 opts + format("-D dcn=%d -D bidx=%d", dcn, bidx));
+        k.create("YUV2RGB_NVx", ocl::imgproc::cvtcolor_oclsrc,
+                 opts + format("-D dcn=%d -D bidx=%d -D uidx=%d", dcn, bidx, uidx));
         break;
     }
     case COLOR_BGR2YCrCb:

modules/imgproc/src/opencl/cvtcolor.cl

@@ -111,6 +111,10 @@ enum
 #define B_COMP w
 #endif
 
+#ifndef uidx
+#define uidx 0
+#endif
+
 #define __CAT(x, y) x##y
 #define CAT(x, y) __CAT(x, y)
 
@@ -297,7 +301,7 @@ __constant int ITUR_BT_601_CVG = 852492;
 __constant int ITUR_BT_601_CVR = 1673527;
 __constant int ITUR_BT_601_SHIFT = 20;
 
-__kernel void YUV2RGB_NV12(__global const uchar* srcptr, int src_step, int src_offset,
+__kernel void YUV2RGB_NVx(__global const uchar* srcptr, int src_step, int src_offset,
                             __global uchar* dstptr, int dst_step, int dt_offset,
                             int rows, int cols)
 {
@@ -321,8 +325,8 @@ __kernel void YUV2RGB_NV12(__global const uchar* srcptr, int src_step, int src_o
                 int Y3 = ysrc[src_step];
                 int Y4 = ysrc[src_step + 1];
 
-                int U  = usrc[0] - 128;
-                int V  = usrc[1] - 128;
+                int U  = usrc[uidx] - 128;
+                int V  = usrc[1 - uidx] - 128;
 
                 int ruv = (1 << (ITUR_BT_601_SHIFT - 1)) + ITUR_BT_601_CVR * V;
                 int guv = (1 << (ITUR_BT_601_SHIFT - 1)) - ITUR_BT_601_CVG * V - ITUR_BT_601_CUG * U;

modules/imgproc/test/ocl/test_color.cpp

@@ -320,7 +320,7 @@ OCL_TEST_P(CvtColor8u32f, Luv2RGBA) { performTest(3, 4, CVTCODE(Luv2RGB), depth
 OCL_TEST_P(CvtColor8u32f, Luv2LBGRA) { performTest(3, 4, CVTCODE(Luv2LBGR), depth == CV_8U ? 1 : 1e-5); }
 OCL_TEST_P(CvtColor8u32f, Luv2LRGBA) { performTest(3, 4, CVTCODE(Luv2LRGB), depth == CV_8U ? 1 : 1e-5); }
 
-// YUV -> RGBA_NV12
+// YUV -> RGBA_NVx
 
 struct CvtColor_YUV420 :
         public CvtColor
@@ -348,6 +348,10 @@ OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) { performTest(1, 4, COLOR_YUV2RGBA_NV
 OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) { performTest(1, 4, COLOR_YUV2BGRA_NV12); }
 OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { performTest(1, 3, COLOR_YUV2RGB_NV12); }
 OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { performTest(1, 3, COLOR_YUV2BGR_NV12); }
+OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV21) { performTest(1, 4, COLOR_YUV2RGBA_NV21); }
+OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV21) { performTest(1, 4, COLOR_YUV2BGRA_NV21); }
+OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV21) { performTest(1, 3, COLOR_YUV2RGB_NV21); }
+OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV21) { performTest(1, 3, COLOR_YUV2BGR_NV21); }
 
 
 OCL_INSTANTIATE_TEST_CASE_P(ImgProc, CvtColor8u,