Merge pull request #2340 from ilya-lavrenov:tapi_lab

modules/imgproc/perf/opencl/perf_color.cpp

@@ -58,7 +58,8 @@ 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_HLS2RGB, COLOR_BGR5652BGR, COLOR_BGR2BGR565, COLOR_RGBA2mRGBA, COLOR_mRGBA2RGBA, COLOR_YUV2RGB_NV12,
+        COLOR_RGB2Lab, COLOR_Lab2BGR)
 
 typedef tuple<Size, tuple<ConversionTypes, int, int> > CvtColorParams;
 typedef TestBaseWithParam<CvtColorParams> CvtColorFixture;
@@ -82,7 +83,9 @@ OCL_PERF_TEST_P(CvtColorFixture, CvtColor, testing::Combine(
                     make_tuple(ConversionTypes(COLOR_BGR2BGR565), 3, 2),
                     make_tuple(ConversionTypes(COLOR_RGBA2mRGBA), 4, 4),
                     make_tuple(ConversionTypes(COLOR_mRGBA2RGBA), 4, 4),
-                    make_tuple(ConversionTypes(COLOR_YUV2RGB_NV12), 1, 3)
+                    make_tuple(ConversionTypes(COLOR_YUV2RGB_NV12), 1, 3),
+                    make_tuple(ConversionTypes(COLOR_RGB2Lab), 3, 3),
+                    make_tuple(ConversionTypes(COLOR_Lab2BGR), 3, 4)
                     )))
 {
     CvtColorParams params = GetParam();

modules/imgproc/src/color.cpp

@@ -1594,8 +1594,11 @@ struct RGB2Lab_b
         static volatile int _3 = 3;
         initLabTabs();
 
-        if(!_coeffs) _coeffs = sRGB2XYZ_D65;
+        if (!_coeffs)
-        if(!_whitept) _whitept = D65;
+            _coeffs = sRGB2XYZ_D65;
+        if (!_whitept)
+            _whitept = D65;
+
         float scale[] =
         {
             (1 << lab_shift)/_whitept[0],
@@ -1699,10 +1702,6 @@ struct RGB2Lab_f
             float G = clip(src[1]);
             float B = clip(src[2]);
 
-//            CV_Assert(R >= 0.0f && R <= 1.0f);
-//            CV_Assert(G >= 0.0f && G <= 1.0f);
-//            CV_Assert(B >= 0.0f && B <= 1.0f);
-
             if (gammaTab)
             {
                 R = splineInterpolate(R * gscale, gammaTab, GAMMA_TAB_SIZE);
@@ -1738,7 +1737,7 @@ struct Lab2RGB_f
 
     Lab2RGB_f( int _dstcn, int blueIdx, const float* _coeffs,
               const float* _whitept, bool _srgb )
-    : dstcn(_dstcn), srgb(_srgb), blueInd(blueIdx)
+    : dstcn(_dstcn), srgb(_srgb)
     {
         initLabTabs();
 
@@ -1796,13 +1795,12 @@ struct Lab2RGB_f
 
 
             float x = fxz[0], z = fxz[1];
-            float ro = clip(C0 * x + C1 * y + C2 * z);
+            float ro = C0 * x + C1 * y + C2 * z;
-            float go = clip(C3 * x + C4 * y + C5 * z);
+            float go = C3 * x + C4 * y + C5 * z;
-            float bo = clip(C6 * x + C7 * y + C8 * z);
+            float bo = C6 * x + C7 * y + C8 * z;
-
+            ro = clip(ro);
-//            CV_Assert(ro >= 0.0f && ro <= 1.0f);
+            go = clip(go);
-//            CV_Assert(go >= 0.0f && go <= 1.0f);
+            bo = clip(bo);
-//            CV_Assert(bo >= 0.0f && bo <= 1.0f);
 
             if (gammaTab)
             {
@@ -1820,7 +1818,6 @@ struct Lab2RGB_f
     int dstcn;
     float coeffs[9];
     bool srgb;
-    int blueInd;
 };
 
 #undef clip
@@ -2700,7 +2697,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
     size_t globalsize[] = { src.cols, src.rows };
     ocl::Kernel k;
 
-    if(depth != CV_8U && depth != CV_16U && depth != CV_32F)
+    if (depth != CV_8U && depth != CV_16U && depth != CV_32F)
         return false;
 
     switch (code)
@@ -3028,8 +3025,162 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
                  format("-D depth=%d -D dcn=4 -D scn=4 -D bidx=3", depth));
         break;
     }
+    case CV_BGR2Lab: case CV_RGB2Lab: case CV_LBGR2Lab: case CV_LRGB2Lab:
+    {
+        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;
+        dcn = 3;
+
+        k.create("BGR2Lab", ocl::imgproc::cvtcolor_oclsrc,
+                 format("-D depth=%d -D dcn=3 -D scn=%d -D bidx=%d%s",
+                        depth, scn, bidx, srgb ? " -D SRGB" : ""));
+        if (k.empty())
+            return false;
+
+        initLabTabs();
+
+        _dst.create(dstSz, CV_MAKETYPE(depth, dcn));
+        dst = _dst.getUMat();
+
+        ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
+                dstarg = ocl::KernelArg::WriteOnly(dst);
+
+        if (depth == CV_8U)
+        {
+            static UMat usRGBGammaTab, ulinearGammaTab, uLabCbrtTab, ucoeffs;
+
+            if (srgb && usRGBGammaTab.empty())
+                Mat(1, 256, CV_16UC1, sRGBGammaTab_b).copyTo(usRGBGammaTab);
+            else if (ulinearGammaTab.empty())
+                Mat(1, 256, CV_16UC1, linearGammaTab_b).copyTo(ulinearGammaTab);
+            if (uLabCbrtTab.empty())
+                Mat(1, LAB_CBRT_TAB_SIZE_B, CV_16UC1, LabCbrtTab_b).copyTo(uLabCbrtTab);
+
+            {
+                int coeffs[9];
+                const float * const _coeffs = sRGB2XYZ_D65, * const _whitept = D65;
+                const float scale[] =
+                {
+                    (1 << lab_shift)/_whitept[0],
+                    (float)(1 << lab_shift),
+                    (1 << lab_shift)/_whitept[2]
+                };
+
+                for (int i = 0; i < 3; i++ )
+                {
+                    coeffs[i*3+(bidx^2)] = cvRound(_coeffs[i*3]*scale[i]);
+                    coeffs[i*3+1] = cvRound(_coeffs[i*3+1]*scale[i]);
+                    coeffs[i*3+bidx] = cvRound(_coeffs[i*3+2]*scale[i]);
+
+                    CV_Assert( coeffs[i] >= 0 && coeffs[i*3+1] >= 0 && coeffs[i*3+2] >= 0 &&
+                              coeffs[i*3] + coeffs[i*3+1] + coeffs[i*3+2] < 2*(1 << lab_shift) );
+                }
+                Mat(1, 9, CV_32SC1, coeffs).copyTo(ucoeffs);
+            }
+
+            const int Lscale = (116*255+50)/100;
+            const int Lshift = -((16*255*(1 << lab_shift2) + 50)/100);
+
+            k.args(srcarg, dstarg,
+                   ocl::KernelArg::PtrReadOnly(srgb ? usRGBGammaTab : ulinearGammaTab),
+                   ocl::KernelArg::PtrReadOnly(uLabCbrtTab), ocl::KernelArg::PtrReadOnly(ucoeffs),
+                   Lscale, Lshift);
+        }
+        else
+        {
+            static UMat usRGBGammaTab, ucoeffs;
+
+            if (srgb && usRGBGammaTab.empty())
+                Mat(1, GAMMA_TAB_SIZE * 4, CV_32FC1, sRGBGammaTab).copyTo(usRGBGammaTab);
+
+            {
+                float coeffs[9];
+                const float * const _coeffs = sRGB2XYZ_D65, * const _whitept = D65;
+                float scale[] = { 1.0f / _whitept[0], 1.0f, 1.0f / _whitept[2] };
+
+                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];
+
+                    CV_Assert( coeffs[j] >= 0 && coeffs[j + 1] >= 0 && coeffs[j + 2] >= 0 &&
+                               coeffs[j] + coeffs[j + 1] + coeffs[j + 2] < 1.5f*LabCbrtTabScale );
+                }
+
+                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);
+            else
+                k.args(srcarg, dstarg, ucoeffsarg, _1_3, _a);
+        }
+
+        return k.run(dims, globalsize, NULL, false);
+    }
+    case CV_Lab2BGR: case CV_Lab2RGB: case CV_Lab2LBGR: case CV_Lab2LRGB:
+    {
+        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;
+
+        k.create("Lab2BGR", ocl::imgproc::cvtcolor_oclsrc,
+                 format("-D depth=%d -D dcn=%d -D scn=3 -D bidx=%d%s",
+                        depth, dcn, bidx, srgb ? " -D SRGB" : ""));
+        if (k.empty())
+            return false;
+
+        initLabTabs();
+        static UMat ucoeffs, usRGBInvGammaTab;
+
+        if (srgb && usRGBInvGammaTab.empty())
+            Mat(1, GAMMA_TAB_SIZE*4, CV_32FC1, sRGBInvGammaTab).copyTo(usRGBInvGammaTab);
+
+        {
+            float coeffs[9];
+            const float * const _coeffs = XYZ2sRGB_D65, * const _whitept = D65;
+
+            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];
+            }
+
+            Mat(1, 9, CV_32FC1, coeffs).copyTo(ucoeffs);
+        }
+
+        _dst.create(sz, CV_MAKETYPE(depth, dcn));
+        dst = _dst.getUMat();
+
+        float lThresh = 0.008856f * 903.3f;
+        float fThresh = 7.787f * 0.008856f + 16.0f / 116.0f;
+
+        ocl::KernelArg srcarg = ocl::KernelArg::ReadOnlyNoSize(src),
+                dstarg = ocl::KernelArg::WriteOnly(dst),
+                coeffsarg = ocl::KernelArg::PtrReadOnly(ucoeffs);
+
+        if (srgb)
+            k.args(srcarg, dstarg, ocl::KernelArg::PtrReadOnly(usRGBInvGammaTab),
+                   coeffsarg, lThresh, fThresh);
+        else
+            k.args(srcarg, dstarg, coeffsarg, lThresh, fThresh);
+
+        return k.run(dims, globalsize, NULL, false);
+    }
     default:
-        ;
+        break;
     }
 
     if( !k.empty() )
@@ -3037,7 +3188,7 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
         _dst.create(dstSz, CV_MAKETYPE(depth, dcn));
         dst = _dst.getUMat();
         k.args(ocl::KernelArg::ReadOnlyNoSize(src), ocl::KernelArg::WriteOnly(dst));
-        ok = k.run(dims, globalsize, 0, false);
+        ok = k.run(dims, globalsize, NULL, false);
     }
     return ok;
 }
@@ -3771,9 +3922,9 @@ void cv::cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
                 dst = _dst.getMat();
 
                 if( depth == CV_8U )
-                {
                     CvtColorLoop(src, dst, RGBA2mRGBA<uchar>());
-                } else {
+                else
+                {
                     CV_Error( CV_StsBadArg, "Unsupported image depth" );
                 }
             }
@@ -3787,9 +3938,9 @@ void cv::cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
                 dst = _dst.getMat();
 
                 if( depth == CV_8U )
-                {
                     CvtColorLoop(src, dst, mRGBA2RGBA<uchar>());
-                } else {
+                else
+                {
                     CV_Error( CV_StsBadArg, "Unsupported image depth" );
                 }
             }

modules/imgproc/src/opencl/cvtcolor.cl

@@ -46,10 +46,6 @@
 
 /**************************************PUBLICFUNC*************************************/
 
-#if defined (DOUBLE_SUPPORT)
-#pragma OPENCL EXTENSION cl_khr_fp64:enable
-#endif
-
 #if depth == 0
     #define DATA_TYPE uchar
     #define MAX_NUM  255
@@ -1065,3 +1061,234 @@ __kernel void mRGBA2RGBA(__global const uchar* src, int src_step, int src_offset
 }
 
 #endif
+
+/////////////////////////////////// [l|s]RGB <-> Lab ///////////////////////////
+
+#define lab_shift xyz_shift
+#define gamma_shift 3
+#define lab_shift2 (lab_shift + gamma_shift)
+#define GAMMA_TAB_SIZE 1024
+#define GammaTabScale (float)GAMMA_TAB_SIZE
+
+inline float splineInterpolate(float x, __global const float * tab, int n)
+{
+    int ix = clamp(convert_int_sat_rtn(x), 0, n-1);
+    x -= ix;
+    tab += ix*4;
+    return ((tab[3]*x + tab[2])*x + tab[1])*x + tab[0];
+}
+
+#ifdef DEPTH_0
+
+__kernel void BGR2Lab(__global const uchar * src, int src_step, int src_offset,
+                      __global uchar * dst, int dst_step, int dst_offset, int rows, int cols,
+                      __global const ushort * gammaTab, __global ushort * LabCbrtTab_b,
+                      __constant int * coeffs, int Lscale, int Lshift)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (y < rows && x < cols)
+    {
+        int src_idx = mad24(y, src_step, src_offset + x * scnbytes);
+        int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes);
+
+        src += src_idx;
+        dst += dst_idx;
+
+        int C0 = coeffs[0], C1 = coeffs[1], C2 = coeffs[2],
+            C3 = coeffs[3], C4 = coeffs[4], C5 = coeffs[5],
+            C6 = coeffs[6], C7 = coeffs[7], C8 = coeffs[8];
+
+        int R = gammaTab[src[0]], G = gammaTab[src[1]], B = gammaTab[src[2]];
+        int fX = LabCbrtTab_b[CV_DESCALE(R*C0 + G*C1 + B*C2, lab_shift)];
+        int fY = LabCbrtTab_b[CV_DESCALE(R*C3 + G*C4 + B*C5, lab_shift)];
+        int fZ = LabCbrtTab_b[CV_DESCALE(R*C6 + G*C7 + B*C8, lab_shift)];
+
+        int L = CV_DESCALE( Lscale*fY + Lshift, lab_shift2 );
+        int a = CV_DESCALE( 500*(fX - fY) + 128*(1 << lab_shift2), lab_shift2 );
+        int b = CV_DESCALE( 200*(fY - fZ) + 128*(1 << lab_shift2), lab_shift2 );
+
+        dst[0] = SAT_CAST(L);
+        dst[1] = SAT_CAST(a);
+        dst[2] = SAT_CAST(b);
+    }
+}
+
+#elif defined DEPTH_5
+
+__kernel void BGR2Lab(__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 _1_3, float _a)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (y < rows && x < cols)
+    {
+        int src_idx = mad24(y, src_step, src_offset + x * scnbytes);
+        int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes);
+
+        __global const float * src = (__global const float *)(srcptr + src_idx);
+        __global float * dst = (__global float *)(dstptr + dst_idx);
+
+        float C0 = coeffs[0], C1 = coeffs[1], C2 = coeffs[2],
+              C3 = coeffs[3], C4 = coeffs[4], C5 = coeffs[5],
+              C6 = coeffs[6], C7 = coeffs[7], C8 = coeffs[8];
+
+        float R = clamp(src[0], 0.0f, 1.0f);
+        float G = clamp(src[1], 0.0f, 1.0f);
+        float B = clamp(src[2], 0.0f, 1.0f);
+
+#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*C0 + G*C1 + B*C2;
+        float Y = R*C3 + G*C4 + B*C5;
+        float Z = R*C6 + G*C7 + B*C8;
+
+        float FX = X > 0.008856f ? pow(X, _1_3) : (7.787f * X + _a);
+        float FY = Y > 0.008856f ? pow(Y, _1_3) : (7.787f * Y + _a);
+        float FZ = Z > 0.008856f ? pow(Z, _1_3) : (7.787f * Z + _a);
+
+        float L = Y > 0.008856f ? (116.f * FY - 16.f) : (903.3f * Y);
+        float a = 500.f * (FX - FY);
+        float b = 200.f * (FY - FZ);
+
+        dst[0] = L;
+        dst[1] = a;
+        dst[2] = b;
+    }
+}
+
+#endif
+
+inline void Lab2BGR_f(const float * srcbuf, float * dstbuf,
+#ifdef SRGB
+                      __global const float * gammaTab,
+#endif
+                      __constant float * coeffs, float lThresh, float fThresh)
+{
+    float li = srcbuf[0], ai = srcbuf[1], bi = srcbuf[2];
+
+    float C0 = coeffs[0], C1 = coeffs[1], C2 = coeffs[2],
+          C3 = coeffs[3], C4 = coeffs[4], C5 = coeffs[5],
+          C6 = coeffs[6], C7 = coeffs[7], C8 = coeffs[8];
+
+    float y, fy;
+    if (li <= lThresh)
+    {
+        y = li / 903.3f;
+        fy = 7.787f * y + 16.0f / 116.0f;
+    }
+    else
+    {
+        fy = (li + 16.0f) / 116.0f;
+        y = fy * fy * fy;
+    }
+
+    float fxz[] = { ai / 500.0f + fy, fy - bi / 200.0f };
+
+    for (int j = 0; j < 2; j++)
+        if (fxz[j] <= fThresh)
+            fxz[j] = (fxz[j] - 16.0f / 116.0f) / 7.787f;
+        else
+            fxz[j] = fxz[j] * fxz[j] * fxz[j];
+
+    float x = fxz[0], z = fxz[1];
+    float ro = clamp(C0 * x + C1 * y + C2 * z, 0.0f, 1.0f);
+    float go = clamp(C3 * x + C4 * y + C5 * z, 0.0f, 1.0f);
+    float bo = clamp(C6 * x + C7 * y + C8 * z, 0.0f, 1.0f);
+
+#ifdef SRGB
+    ro = splineInterpolate(ro * GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+    go = splineInterpolate(go * GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+    bo = splineInterpolate(bo * GammaTabScale, gammaTab, GAMMA_TAB_SIZE);
+#endif
+
+    dstbuf[0] = ro, dstbuf[1] = go, dstbuf[2] = bo;
+}
+
+#ifdef DEPTH_0
+
+__kernel void Lab2BGR(__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 lThresh, float fThresh)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (y < rows && x < cols)
+    {
+        int src_idx = mad24(y, src_step, src_offset + x * scnbytes);
+        int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes);
+
+        src += src_idx;
+        dst += dst_idx;
+
+        float srcbuf[3], dstbuf[3];
+        srcbuf[0] = src[0]*(100.f/255.f);
+        srcbuf[1] = convert_float(src[1] - 128);
+        srcbuf[2] = convert_float(src[2] - 128);
+
+        Lab2BGR_f(&srcbuf[0], &dstbuf[0],
+#ifdef SRGB
+            gammaTab,
+#endif
+            coeffs, lThresh, fThresh);
+
+        dst[0] = SAT_CAST(dstbuf[0] * 255.0f);
+        dst[1] = SAT_CAST(dstbuf[1] * 255.0f);
+        dst[2] = SAT_CAST(dstbuf[2] * 255.0f);
+#if dcn == 4
+        dst[3] = MAX_NUM;
+#endif
+    }
+}
+
+#elif defined DEPTH_5
+
+__kernel void Lab2BGR(__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 lThresh, float fThresh)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (y < rows && x < cols)
+    {
+        int src_idx = mad24(y, src_step, src_offset + x * scnbytes);
+        int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes);
+
+        __global const float * src = (__global const float *)(srcptr + src_idx);
+        __global float * dst = (__global float *)(dstptr + dst_idx);
+
+        float srcbuf[3], dstbuf[3];
+        srcbuf[0] = src[0], srcbuf[1] = src[1], srcbuf[2] = src[2];
+
+        Lab2BGR_f(&srcbuf[0], &dstbuf[0],
+#ifdef SRGB
+            gammaTab,
+#endif
+            coeffs, lThresh, fThresh);
+
+        dst[0] = dstbuf[0], dst[1] = dstbuf[1], dst[2] = dstbuf[2];
+#if dcn == 4
+        dst[3] = MAX_NUM;
+#endif
+    }
+}
+
+#endif

modules/imgproc/test/ocl/test_color.cpp

@@ -248,6 +248,26 @@ OCL_TEST_P(CvtColor8u, GRAY2BGR555) { performTest(1, 2, CVTCODE(GRAY2BGR555)); }
 OCL_TEST_P(CvtColor8u, RGBA2mRGBA) { performTest(4, 4, CVTCODE(RGBA2mRGBA)); }
 OCL_TEST_P(CvtColor8u, mRGBA2RGBA) { performTest(4, 4, CVTCODE(mRGBA2RGBA)); }
 
+// RGB <-> Lab
+
+OCL_TEST_P(CvtColor8u32f, BGR2Lab) { performTest(3, 3, CVTCODE(BGR2Lab)); }
+OCL_TEST_P(CvtColor8u32f, RGB2Lab) { performTest(3, 3, CVTCODE(RGB2Lab)); }
+OCL_TEST_P(CvtColor8u32f, LBGR2Lab) { performTest(3, 3, CVTCODE(LBGR2Lab)); }
+OCL_TEST_P(CvtColor8u32f, LRGB2Lab) { performTest(3, 3, CVTCODE(LRGB2Lab)); }
+OCL_TEST_P(CvtColor8u32f, BGRA2Lab) { performTest(4, 3, CVTCODE(BGR2Lab)); }
+OCL_TEST_P(CvtColor8u32f, RGBA2Lab) { performTest(4, 3, CVTCODE(RGB2Lab)); }
+OCL_TEST_P(CvtColor8u32f, LBGRA2Lab) { performTest(4, 3, CVTCODE(LBGR2Lab)); }
+OCL_TEST_P(CvtColor8u32f, LRGBA2Lab) { performTest(4, 3, CVTCODE(LRGB2Lab)); }
+
+OCL_TEST_P(CvtColor8u32f, Lab2BGR) { performTest(3, 3, CVTCODE(Lab2BGR), depth == CV_8U ? 1 : 1e-5); }
+OCL_TEST_P(CvtColor8u32f, Lab2RGB) { performTest(3, 3, CVTCODE(Lab2RGB), depth == CV_8U ? 1 : 1e-5); }
+OCL_TEST_P(CvtColor8u32f, Lab2LBGR) { performTest(3, 3, CVTCODE(Lab2LBGR), depth == CV_8U ? 1 : 1e-5); }
+OCL_TEST_P(CvtColor8u32f, Lab2LRGB) { performTest(3, 3, CVTCODE(Lab2LRGB), depth == CV_8U ? 1 : 1e-5); }
+OCL_TEST_P(CvtColor8u32f, Lab2BGRA) { performTest(3, 4, CVTCODE(Lab2BGR), depth == CV_8U ? 1 : 1e-5); }
+OCL_TEST_P(CvtColor8u32f, Lab2RGBA) { performTest(3, 4, CVTCODE(Lab2RGB), depth == CV_8U ? 1 : 1e-5); }
+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); }
+
 // YUV -> RGBA_NV12
 
 struct CvtColor_YUV420 :