Merge pull request #2700 from ilya-lavrenov:tapi_cvtColor

modules/imgproc/perf/opencl/perf_color.cpp

@@ -59,7 +59,7 @@ using std::tr1::make_tuple;
 CV_ENUM(ConversionTypes, COLOR_RGB2GRAY, COLOR_RGB2BGR, COLOR_RGB2YUV, COLOR_YUV2RGB, COLOR_RGB2YCrCb,
         COLOR_YCrCb2RGB, COLOR_RGB2XYZ, COLOR_XYZ2RGB, COLOR_RGB2HSV, COLOR_HSV2RGB, COLOR_RGB2HLS,
         COLOR_HLS2RGB, COLOR_BGR5652BGR, COLOR_BGR2BGR565, COLOR_RGBA2mRGBA, COLOR_mRGBA2RGBA, COLOR_YUV2RGB_NV12,
-        COLOR_RGB2Lab, COLOR_Lab2BGR)
+        COLOR_RGB2Lab, COLOR_Lab2BGR, COLOR_RGB2Luv, COLOR_Luv2LBGR)
 
 typedef tuple<Size, tuple<ConversionTypes, int, int> > CvtColorParams;
 typedef TestBaseWithParam<CvtColorParams> CvtColorFixture;
@@ -85,7 +85,9 @@ OCL_PERF_TEST_P(CvtColorFixture, CvtColor, testing::Combine(
                     make_tuple(ConversionTypes(COLOR_mRGBA2RGBA), 4, 4),
                     make_tuple(ConversionTypes(COLOR_YUV2RGB_NV12), 1, 3),
                     make_tuple(ConversionTypes(COLOR_RGB2Lab), 3, 3),
-                    make_tuple(ConversionTypes(COLOR_Lab2BGR), 3, 4)
+                    make_tuple(ConversionTypes(COLOR_Lab2BGR), 3, 4),
+                    make_tuple(ConversionTypes(COLOR_RGB2Luv), 3, 3),
+                    make_tuple(ConversionTypes(COLOR_Luv2LBGR), 3, 4)
                     )))
 {
     CvtColorParams params = GetParam();

modules/imgproc/src/color.cpp

@@ -2749,12 +2749,13 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
 
     ocl::Device dev = ocl::Device::getDefault();
     int pxPerWIy = 1;
-    if (dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU))
-    {
+    if (dev.isIntel() && (dev.type() & ocl::Device::TYPE_GPU) &&
+            !(code == CV_BGR2Luv || code == CV_RGB2Luv || code == CV_LBGR2Luv || code == CV_LRGB2Luv ||
+              code == CV_Luv2BGR || code == CV_Luv2RGB || code == CV_Luv2LBGR || code == CV_Luv2LRGB))
         pxPerWIy = 4;
-    }
+
     globalsize[1] = DIVUP(globalsize[1], pxPerWIy);
-    opts +=  format("-D PIX_PER_WI_Y=%d ", pxPerWIy);
+    opts += format("-D PIX_PER_WI_Y=%d ", pxPerWIy);
 
     switch (code)
     {
@@ -3084,16 +3085,20 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
         break;
     }
     case CV_BGR2Lab: case CV_RGB2Lab: case CV_LBGR2Lab: case CV_LRGB2Lab:
+    case CV_BGR2Luv: case CV_RGB2Luv: case CV_LBGR2Luv: case CV_LRGB2Luv:
     {
         CV_Assert( (scn == 3 || scn == 4) && (depth == CV_8U || depth == CV_32F) );
 
-        bidx = code == CV_BGR2Lab || code == CV_LBGR2Lab ? 0 : 2;
-        bool srgb = code == CV_BGR2Lab || code == CV_RGB2Lab;
+        bidx = code == CV_BGR2Lab || code == CV_LBGR2Lab || code == CV_BGR2Luv || code == CV_LBGR2Luv ? 0 : 2;
+        bool srgb = code == CV_BGR2Lab || code == CV_RGB2Lab || code == CV_RGB2Luv || code == CV_BGR2Luv;
+        bool lab = code == CV_BGR2Lab || code == CV_RGB2Lab || code == CV_LBGR2Lab || code == CV_LRGB2Lab;
+        float un, vn;
         dcn = 3;
 
-        k.create("BGR2Lab", ocl::imgproc::cvtcolor_oclsrc,
-                 opts + format("-D dcn=3 -D bidx=%d%s",
-                               bidx, srgb ? " -D SRGB" : ""));
+        k.create(format("BGR2%s", lab ? "Lab" : "Luv").c_str(),
+                 ocl::imgproc::cvtcolor_oclsrc,
+                 opts + format("-D dcn=%d -D bidx=%d%s",
+                               dcn, bidx, srgb ? " -D SRGB" : ""));
         if (k.empty())
             return false;
 
@@ -3105,7 +3110,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
         ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
                 dstarg = ocl::KernelArg::WriteOnly(dst);
 
-        if (depth == CV_8U)
+        if (depth == CV_8U && lab)
         {
             static UMat usRGBGammaTab, ulinearGammaTab, uLabCbrtTab, ucoeffs;
 
@@ -3148,10 +3153,12 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
         }
         else
         {
-            static UMat usRGBGammaTab, ucoeffs;
+            static UMat usRGBGammaTab, ucoeffs, uLabCbrtTab;
 
             if (srgb && usRGBGammaTab.empty())
                 Mat(1, GAMMA_TAB_SIZE * 4, CV_32FC1, sRGBGammaTab).copyTo(usRGBGammaTab);
+            if (!lab && uLabCbrtTab.empty())
+                Mat(1, LAB_CBRT_TAB_SIZE * 4, CV_32FC1, LabCbrtTab).copyTo(uLabCbrtTab);
 
             {
                 float coeffs[9];
@@ -3161,39 +3168,59 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
                 for (int i = 0; i < 3; i++)
                 {
                     int j = i * 3;
-                    coeffs[j + (bidx ^ 2)] = _coeffs[j] * scale[i];
-                    coeffs[j + 1] = _coeffs[j + 1] * scale[i];
-                    coeffs[j + bidx] = _coeffs[j + 2] * scale[i];
+                    coeffs[j + (bidx ^ 2)] = _coeffs[j] * (lab ? scale[i] : 1);
+                    coeffs[j + 1] = _coeffs[j + 1] * (lab ? scale[i] : 1);
+                    coeffs[j + bidx] = _coeffs[j + 2] * (lab ? scale[i] : 1);
 
                     CV_Assert( coeffs[j] >= 0 && coeffs[j + 1] >= 0 && coeffs[j + 2] >= 0 &&
-                               coeffs[j] + coeffs[j + 1] + coeffs[j + 2] < 1.5f*LabCbrtTabScale );
+                               coeffs[j] + coeffs[j + 1] + coeffs[j + 2] < 1.5f*(lab ? LabCbrtTabScale : 1) );
                 }
 
+                float d = 1.f/(_whitept[0] + _whitept[1]*15 + _whitept[2]*3);
+                un = 13*4*_whitept[0]*d;
+                vn = 13*9*_whitept[1]*d;
+
                 Mat(1, 9, CV_32FC1, coeffs).copyTo(ucoeffs);
             }
 
             float _1_3 = 1.0f / 3.0f, _a = 16.0f / 116.0f;
             ocl::KernelArg ucoeffsarg = ocl::KernelArg::PtrReadOnly(ucoeffs);
 
-            if (srgb)
-                k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBGammaTab),
-                       ucoeffsarg, _1_3, _a);
+            if (lab)
+            {
+                if (srgb)
+                    k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBGammaTab),
+                           ucoeffsarg, _1_3, _a);
+                else
+                    k.args(srcarg, dstarg, ucoeffsarg, _1_3, _a);
+            }
             else
-                k.args(srcarg, dstarg, ucoeffsarg, _1_3, _a);
+            {
+                ocl::KernelArg LabCbrtTabarg = ocl::KernelArg::PtrReadOnly(uLabCbrtTab);
+                if (srgb)
+                    k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBGammaTab),
+                           LabCbrtTabarg, ucoeffsarg, un, vn);
+                else
+                    k.args(srcarg, dstarg, LabCbrtTabarg, ucoeffsarg, un, vn);
+            }
         }
 
         return k.run(dims, globalsize, NULL, false);
     }
     case CV_Lab2BGR: case CV_Lab2RGB: case CV_Lab2LBGR: case CV_Lab2LRGB:
+    case CV_Luv2BGR: case CV_Luv2RGB: case CV_Luv2LBGR: case CV_Luv2LRGB:
     {
         if( dcn <= 0 )
             dcn = 3;
         CV_Assert( scn == 3 && (dcn == 3 || dcn == 4) && (depth == CV_8U || depth == CV_32F) );
 
-        bidx = code == CV_Lab2BGR || code == CV_Lab2LBGR ? 0 : 2;
-        bool srgb = code == CV_Lab2BGR || code == CV_Lab2RGB;
+        bidx = code == CV_Lab2BGR || code == CV_Lab2LBGR || code == CV_Luv2BGR || code == CV_Luv2LBGR ? 0 : 2;
+        bool srgb = code == CV_Lab2BGR || code == CV_Lab2RGB || code == CV_Luv2BGR || code == CV_Luv2RGB;
+        bool lab = code == CV_Lab2BGR || code == CV_Lab2RGB || code == CV_Lab2LBGR || code == CV_Lab2LRGB;
+        float un, vn;
 
-        k.create("Lab2BGR", ocl::imgproc::cvtcolor_oclsrc,
+        k.create(format("%s2BGR", lab ? "Lab" : "Luv").c_str(),
+                 ocl::imgproc::cvtcolor_oclsrc,
                  opts + format("-D dcn=%d -D bidx=%d%s",
                                dcn, bidx, srgb ? " -D SRGB" : ""));
         if (k.empty())
@@ -3211,11 +3238,15 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
 
             for( int i = 0; i < 3; i++ )
             {
-                coeffs[i+(bidx^2)*3] = _coeffs[i]*_whitept[i];
-                coeffs[i+3] = _coeffs[i+3]*_whitept[i];
-                coeffs[i+bidx*3] = _coeffs[i+6]*_whitept[i];
+                coeffs[i+(bidx^2)*3] = _coeffs[i] * (lab ? _whitept[i] : 1);
+                coeffs[i+3] = _coeffs[i+3] * (lab ? _whitept[i] : 1);
+                coeffs[i+bidx*3] = _coeffs[i+6] * (lab ? _whitept[i] : 1);
             }
 
+            float d = 1.f/(_whitept[0] + _whitept[1]*15 + _whitept[2]*3);
+            un = 4*_whitept[0]*d;
+            vn = 9*_whitept[1]*d;
+
             Mat(1, 9, CV_32FC1, coeffs).copyTo(ucoeffs);
         }
 
@@ -3229,11 +3260,22 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
                 dstarg = ocl::KernelArg::WriteOnly(dst),
                 coeffsarg = ocl::KernelArg::PtrReadOnly(ucoeffs);
 
-        if (srgb)
-            k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBInvGammaTab),
-                   coeffsarg, lThresh, fThresh);
+        if (lab)
+        {
+            if (srgb)
+                k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBInvGammaTab),
+                       coeffsarg, lThresh, fThresh);
+            else
+                k.args(srcarg, dstarg, coeffsarg, lThresh, fThresh);
+        }
         else
-            k.args(srcarg, dstarg, coeffsarg, lThresh, fThresh);
+        {
+            if (srgb)
+                k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBInvGammaTab),
+                       coeffsarg, un, vn);
+            else
+                k.args(srcarg, dstarg, coeffsarg, un, vn);
+        }
 
         return k.run(dims, globalsize, NULL, false);
     }
@@ -3264,7 +3306,7 @@ void cv::cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
     int stype = _src.type();
     int scn = CV_MAT_CN(stype), depth = CV_MAT_DEPTH(stype), bidx;
 
-    CV_OCL_RUN( _src.dims() <= 2 && _dst.isUMat(),
+    CV_OCL_RUN( _src.dims() <= 2 && _dst.isUMat() && !(depth == CV_8U && (code == CV_Luv2BGR || code == CV_Luv2RGB)),
                 ocl_cvtColor(_src, _dst, code, dcn) )
 
     Mat src = _src.getMat(), dst;

modules/imgproc/src/opencl/cvtcolor.cl

@@ -1537,3 +1537,205 @@ __kernel void Lab2BGR(__global const uchar * srcptr, int src_step, int src_offse
 }
 
 #endif
+
+/////////////////////////////////// [l|s]RGB <-> Luv ///////////////////////////
+
+#define LAB_CBRT_TAB_SIZE 1024
+#define LAB_CBRT_TAB_SIZE_B (256*3/2*(1<<gamma_shift))
+
+__constant float LabCbrtTabScale = LAB_CBRT_TAB_SIZE/1.5f;
+
+#ifdef DEPTH_5
+
+__kernel void BGR2Luv(__global const uchar * srcptr, int src_step, int src_offset,
+                      __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols,
+#ifdef SRGB
+                      __global const float * gammaTab,
+#endif
+                      __global const float * LabCbrtTab, __constant float * coeffs, float _un, float _vn)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < cols && y < rows)
+    {
+        int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset));
+        int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
+
+        __global const float * src = (__global const float *)(srcptr + src_idx);
+        __global float * dst = (__global float *)(dstptr + dst_idx);
+
+        float R = src[0], G = src[1], B = src[2];
+
+#ifdef SRGB
+        R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+        G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+        B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+#endif
+        float X = R*coeffs[0] + G*coeffs[1] + B*coeffs[2];
+        float Y = R*coeffs[3] + G*coeffs[4] + B*coeffs[5];
+        float Z = R*coeffs[6] + G*coeffs[7] + B*coeffs[8];
+
+        float L = splineInterpolate(Y*LabCbrtTabScale, LabCbrtTab, LAB_CBRT_TAB_SIZE);
+        L = 116.f*L - 16.f;
+
+        float d = (4*13) / max(X + 15 * Y + 3 * Z, FLT_EPSILON);
+        float u = L*(X*d - _un);
+        float v = L*((9*0.25f)*Y*d - _vn);
+
+        dst[0] = L;
+        dst[1] = u;
+        dst[2] = v;
+    }
+}
+
+#elif defined DEPTH_0
+
+__kernel void BGR2Luv(__global const uchar * src, int src_step, int src_offset,
+                      __global uchar * dst, int dst_step, int dst_offset, int rows, int cols,
+#ifdef SRGB
+                      __global const float * gammaTab,
+#endif
+                      __global const float * LabCbrtTab, __constant float * coeffs, float _un, float _vn)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < cols && y < rows)
+    {
+        int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset));
+        int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
+
+        src += src_idx;
+        dst += dst_idx;
+
+        float scale = 1.0f / 255.0f;
+        float R = src[0]*scale, G = src[1]*scale, B = src[2]*scale;
+
+#ifdef SRGB
+        R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+        G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+        B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+#endif
+        float X = R*coeffs[0] + G*coeffs[1] + B*coeffs[2];
+        float Y = R*coeffs[3] + G*coeffs[4] + B*coeffs[5];
+        float Z = R*coeffs[6] + G*coeffs[7] + B*coeffs[8];
+
+        float L = splineInterpolate(Y*LabCbrtTabScale, LabCbrtTab, LAB_CBRT_TAB_SIZE);
+        L = 116.f*L - 16.f;
+
+        float d = (4*13) / max(X + 15 * Y + 3 * Z, FLT_EPSILON);
+        float u = L*(X*d - _un);
+        float v = L*((9*0.25f)*Y*d - _vn);
+
+        dst[0] = SAT_CAST(L * 2.55f);
+        dst[1] = SAT_CAST(mad(u, 0.72033898305084743f, 96.525423728813564f));
+        dst[2] = SAT_CAST(mad(v, 0.99609375f, 139.453125f));
+    }
+}
+
+#endif
+
+#ifdef DEPTH_5
+
+__kernel void Luv2BGR(__global const uchar * srcptr, int src_step, int src_offset,
+                      __global uchar * dstptr, int dst_step, int dst_offset, int rows, int cols,
+#ifdef SRGB
+                      __global const float * gammaTab,
+#endif
+                      __constant float * coeffs, float _un, float _vn)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < cols && y < rows)
+    {
+        int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset));
+        int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
+
+        __global const float * src = (__global const float *)(srcptr + src_idx);
+        __global float * dst = (__global float *)(dstptr + dst_idx);
+
+        float L = src[0], u = src[1], v = src[2], d, X, Y, Z;
+        Y = (L + 16.f) * (1.f/116.f);
+        Y = Y*Y*Y;
+        d = (1.f/13.f)/L;
+        u = u*d + _un;
+        v = v*d + _vn;
+        float iv = 1.f/v;
+        X = 2.25f * u * Y * iv ;
+        Z = (12 - 3 * u - 20 * v) * Y * 0.25f * iv;
+
+        float R = X*coeffs[0] + Y*coeffs[1] + Z*coeffs[2];
+        float G = X*coeffs[3] + Y*coeffs[4] + Z*coeffs[5];
+        float B = X*coeffs[6] + Y*coeffs[7] + Z*coeffs[8];
+
+#ifdef SRGB
+        R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+        G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+        B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+#endif
+
+        dst[0] = R;
+        dst[1] = G;
+        dst[2] = B;
+#if dcn == 4
+        dst[3] = MAX_NUM;
+#endif
+    }
+}
+
+#elif defined DEPTH_0
+
+__kernel void Luv2BGR(__global const uchar * src, int src_step, int src_offset,
+                      __global uchar * dst, int dst_step, int dst_offset, int rows, int cols,
+#ifdef SRGB
+                      __global const float * gammaTab,
+#endif
+                      __constant float * coeffs, float _un, float _vn)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (x < cols && y < rows)
+    {
+        int src_idx = mad24(y, src_step, mad24(x, scnbytes, src_offset));
+        int dst_idx = mad24(y, dst_step, mad24(x, dcnbytes, dst_offset));
+
+        src += src_idx;
+        dst += dst_idx;
+
+        float d, X, Y, Z;
+        float L = src[0]*(100.f/255.f);
+        float u = (float)(src[1]*1.388235294117647f - 134.f);
+        float v = (float)(src[2]*1.003921568627451f - 140.f);
+        Y = (L + 16.f) * (1.f/116.f);
+        Y = Y*Y*Y;
+        d = (1.f/13.f)/L;
+        u = u*d + _un;
+        v = v*d + _vn;
+        float iv = 1.f/v;
+        X = 2.25f * u * Y * iv ;
+        Z = (12 - 3 * u - 20 * v) * Y * 0.25f * iv;
+
+        float R = X*coeffs[0] + Y*coeffs[1] + Z*coeffs[2];
+        float G = X*coeffs[3] + Y*coeffs[4] + Z*coeffs[5];
+        float B = X*coeffs[6] + Y*coeffs[7] + Z*coeffs[8];
+
+#ifdef SRGB
+        R = splineInterpolate(R*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+        G = splineInterpolate(G*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+        B = splineInterpolate(B*GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+#endif
+
+        dst[0] = SAT_CAST(R * 255.0f);
+        dst[1] = SAT_CAST(G * 255.0f);
+        dst[2] = SAT_CAST(B * 255.0f);
+
+#if dcn == 4
+        dst[3] = MAX_NUM;
+#endif
+    }
+}
+
+#endif

modules/imgproc/test/ocl/test_color.cpp

@@ -156,7 +156,7 @@ OCL_TEST_P(CvtColor, YCrCb2BGRA) { performTest(3, 4, CVTCODE(YCrCb2BGR)); }
 // RGB <-> XYZ
 
 #if IPP_VERSION_X100 > 0
-#define IPP_EPS depth <= CV_32S ? 1 : 4e-5
+#define IPP_EPS depth <= CV_32S ? 1 : 5e-5
 #else
 #define IPP_EPS 1e-3
 #endif
@@ -300,6 +300,26 @@ OCL_TEST_P(CvtColor8u32f, Lab2RGBA) { performTest(3, 4, CVTCODE(Lab2RGB), depth
 OCL_TEST_P(CvtColor8u32f, Lab2LBGRA) { performTest(3, 4, CVTCODE(Lab2LBGR), depth == CV_8U ? 1 : 1e-5); }
 OCL_TEST_P(CvtColor8u32f, Lab2LRGBA) { performTest(3, 4, CVTCODE(Lab2LRGB), depth == CV_8U ? 1 : 1e-5); }
 
+// RGB -> Luv
+
+OCL_TEST_P(CvtColor8u32f, BGR2Luv) { performTest(3, 3, CVTCODE(BGR2Luv), depth == CV_8U ? 1 : 1e-2); }
+OCL_TEST_P(CvtColor8u32f, RGB2Luv) { performTest(3, 3, CVTCODE(RGB2Luv), depth == CV_8U ? 1 : 1e-2); }
+OCL_TEST_P(CvtColor8u32f, LBGR2Luv) { performTest(3, 3, CVTCODE(LBGR2Luv), depth == CV_8U ? 1 : 3e-3); }
+OCL_TEST_P(CvtColor8u32f, LRGB2Luv) { performTest(3, 3, CVTCODE(LRGB2Luv), depth == CV_8U ? 1 : 4e-3); }
+OCL_TEST_P(CvtColor8u32f, BGRA2Luv) { performTest(4, 3, CVTCODE(BGR2Luv), depth == CV_8U ? 1 : 8e-3); }
+OCL_TEST_P(CvtColor8u32f, RGBA2Luv) { performTest(4, 3, CVTCODE(RGB2Luv), depth == CV_8U ? 1 : 7e-3); }
+OCL_TEST_P(CvtColor8u32f, LBGRA2Luv) { performTest(4, 3, CVTCODE(LBGR2Luv), depth == CV_8U ? 1 : 4e-3); }
+OCL_TEST_P(CvtColor8u32f, LRGBA2Luv) { performTest(4, 3, CVTCODE(LRGB2Luv), depth == CV_8U ? 1 : 4e-3); }
+
+OCL_TEST_P(CvtColor8u32f, Luv2BGR) { performTest(3, 3, CVTCODE(Luv2BGR), depth == CV_8U ? 1 : 7e-5); }
+OCL_TEST_P(CvtColor8u32f, Luv2RGB) { performTest(3, 3, CVTCODE(Luv2RGB), depth == CV_8U ? 1 : 7e-5); }
+OCL_TEST_P(CvtColor8u32f, Luv2LBGR) { performTest(3, 3, CVTCODE(Luv2LBGR), depth == CV_8U ? 1 : 1e-5); }
+OCL_TEST_P(CvtColor8u32f, Luv2LRGB) { performTest(3, 3, CVTCODE(Luv2LRGB), depth == CV_8U ? 1 : 1e-5); }
+OCL_TEST_P(CvtColor8u32f, Luv2BGRA) { performTest(3, 4, CVTCODE(Luv2BGR), depth == CV_8U ? 1 : 7e-5); }
+OCL_TEST_P(CvtColor8u32f, Luv2RGBA) { performTest(3, 4, CVTCODE(Luv2RGB), depth == CV_8U ? 1 : 7e-5); }
+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
 
 struct CvtColor_YUV420 :