added RGB[A] <-> BGR[A] conversion to ocl::cvtColor

diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp
index 9688be0..f27366a 100644 (file)
--- a/modules/ocl/src/color.cpp
+++ b/modules/ocl/src/color.cpp
@@ -79,7 +79,7 @@ static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::
 static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName,
                          const oclMat & data = oclMat())
 {
-    std::string build_options = format("-D DEPTH_%d -D channels=%d", src.depth(), dst.channels());
+    std::string build_options = format("-D DEPTH_%d -D dcn=%d", src.depth(), dst.channels());
     int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
     int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
 
@@ -101,6 +101,27 @@ static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::st
     openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str());
 }
 
+static void RGB_caller(const oclMat &src, oclMat &dst, bool reverse)
+{
+    std::string build_options = format("-D DEPTH_%d -D dcn=%d -D scn=%d -D %s", src.depth(),
+                                       dst.channels(), src.channels(), reverse ? "REVERSE" : "ORDER");
+    int src_offset = src.offset / src.elemSize1(), src_step = src.step1();
+    int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1();
+
+    vector<pair<size_t , const void *> > args;
+    args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.cols));
+    args.push_back( make_pair( sizeof(cl_int) , (void *)&dst.rows));
+    args.push_back( make_pair( sizeof(cl_int) , (void *)&src_step));
+    args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_step));
+    args.push_back( make_pair( sizeof(cl_mem) , (void *)&src.data));
+    args.push_back( make_pair( sizeof(cl_mem) , (void *)&dst.data));
+    args.push_back( make_pair( sizeof(cl_int) , (void *)&src_offset ));
+    args.push_back( make_pair( sizeof(cl_int) , (void *)&dst_offset ));
+
+    size_t gt[3] = { dst.cols, dst.rows, 1 }, lt[3] = { 16, 16, 1 };
+    openCLExecuteKernel(src.clCxt, &cvt_color, "RGB", gt, lt, args, -1, -1, build_options.c_str());
+}
+
 static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
 {
     Size sz = src.size();
@@ -110,18 +131,24 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
 
     switch (code)
     {
-        /*
-        case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR:
-        case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA:
-        case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:
-        case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555:
-        case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:
-        case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:
-        */
-    case CV_RGB2GRAY:
-    case CV_BGR2GRAY:
-    case CV_RGBA2GRAY:
-    case CV_BGRA2GRAY:
+    case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR:
+    case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA:
+    {
+        CV_Assert(scn == 3 || scn == 4);
+        dcn = code == CV_BGR2BGRA || code == CV_RGB2BGRA || code == CV_BGRA2RGBA ? 4 : 3;
+        bool reverse = !(code == CV_BGR2BGRA || code == CV_BGRA2BGR);
+        dst.create(sz, CV_MAKE_TYPE(depth, dcn));
+        RGB_caller(src, dst, reverse);
+        break;
+    }
+    /*
+    case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555:
+    case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555:
+    case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:
+    case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:
+    */
+    case CV_RGB2GRAY: case CV_BGR2GRAY:
+    case CV_RGBA2GRAY: case CV_BGRA2GRAY:
     {
         CV_Assert(scn == 3 || scn == 4);
         bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2;
@@ -158,10 +185,8 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn)
         toRGB_caller(src, dst, bidx, "YUV2RGB");
         break;
     }
-    case CV_YUV2RGB_NV12:
-    case CV_YUV2BGR_NV12:
-    case CV_YUV2RGBA_NV12:
-    case CV_YUV2BGRA_NV12:
+    case CV_YUV2RGB_NV12: case CV_YUV2BGR_NV12:
+    case CV_YUV2RGBA_NV12: case CV_YUV2BGRA_NV12:
     {
         CV_Assert(scn == 1);
         CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U );

diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl
index 0f44897..916d44b 100644 (file)
--- a/modules/ocl/src/opencl/cvt_color.cl
+++ b/modules/ocl/src/opencl/cvt_color.cl
@@ -119,8 +119,8 @@ __kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int bidx,
         dst[dst_idx] = val;
         dst[dst_idx + 1] = val;
         dst[dst_idx + 2] = val;
-#if channels == 4
-            dst[dst_idx + 3] = MAX_NUM;
+#if dcn == 4
+        dst[dst_idx + 3] = MAX_NUM;
 #endif
     }
 }
@@ -195,8 +195,8 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step,
         dst[dst_idx + bidx] = SAT_CAST( b );
         dst[dst_idx + 1]      = SAT_CAST( g );
         dst[dst_idx + (bidx^2)]   = SAT_CAST( r );
-#if channels == 4
-            dst[dst_idx + 3] = MAX_NUM;
+#if dcn == 4
+        dst[dst_idx + 3] = MAX_NUM;
 #endif
     }
 }
@@ -332,8 +332,8 @@ __kernel void YCrCb2RGB(int cols, int rows, int src_step, int dst_step,
         dst[dst_idx + (bidx^2)] = SAT_CAST(r);
         dst[dst_idx + 1] = SAT_CAST(g);
         dst[dst_idx + bidx] = SAT_CAST(b);
-#if channels == 4
-            dst[dst_idx + 3] = MAX_NUM;
+#if dcn == 4
+        dst[dst_idx + 3] = MAX_NUM;
 #endif
     }
 }
@@ -397,8 +397,43 @@ __kernel void XYZ2RGB(int cols, int rows, int src_step, int dst_step,
         dst[dst_idx] = SAT_CAST(b);
         dst[dst_idx + 1] = SAT_CAST(g);
         dst[dst_idx + 2] = SAT_CAST(r);
-#if channels == 4
-            dst[dst_idx + 3] = MAX_NUM;
+#if dcn == 4
+        dst[dst_idx + 3] = MAX_NUM;
+#endif
+    }
+}
+
+///////////////////////////////////// RGB[A] <-> BGR[A] //////////////////////////////////////
+
+__kernel void RGB(int cols, int rows, int src_step, int dst_step,
+                  __global const DATA_TYPE * src, __global DATA_TYPE * dst,
+                  int src_offset, int dst_offset)
+{
+    int x = get_global_id(0);
+    int y = get_global_id(1);
+
+    if (y < rows && x < cols)
+    {
+        x <<= 2;
+        int src_idx = mad24(y, src_step, src_offset + x);
+        int dst_idx = mad24(y, dst_step, dst_offset + x);
+
+#ifdef REVERSE
+        dst[dst_idx] = src[src_idx + 2];
+        dst[dst_idx + 1] = src[src_idx + 1];
+        dst[dst_idx + 2] = src[src_idx];
+#elif defined ORDER
+        dst[dst_idx] = src[src_idx];
+        dst[dst_idx + 1] = src[src_idx + 1];
+        dst[dst_idx + 2] = src[src_idx + 2];
+#endif
+
+#if dcn == 4
+#if scn == 3
+        dst[dst_idx + 3] = MAX_NUM;
+#else
+        dst[dst_idx + 3] = src[src_idx + 3];
+#endif
 #endif
     }
 }

diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp
index 977fd20..e982684 100644 (file)
--- a/modules/ocl/test/test_color.cpp
+++ b/modules/ocl/test/test_color.cpp
@@ -121,149 +121,65 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool)
 
 #define CVTCODE(name) COLOR_ ## name
 
-// RGB <-> Gray
-
-OCL_TEST_P(CvtColor, RGB2GRAY)
-{
-    doTest(3, 1, CVTCODE(RGB2GRAY));
-}
-OCL_TEST_P(CvtColor, GRAY2RGB)
-{
-    doTest(1, 3, CVTCODE(GRAY2RGB));
-}
+// RGB[A] <-> BGR[A]
+
+OCL_TEST_P(CvtColor, BGR2BGRA) { doTest(3, 4, CVTCODE(BGR2BGRA)); }
+OCL_TEST_P(CvtColor, RGB2RGBA) { doTest(3, 4, CVTCODE(BGR2BGRA)); }
+OCL_TEST_P(CvtColor, BGRA2BGR) { doTest(4, 3, CVTCODE(BGRA2BGR)); }
+OCL_TEST_P(CvtColor, RGBA2RGB) { doTest(4, 3, CVTCODE(BGRA2BGR)); }
+OCL_TEST_P(CvtColor, BGR2RGBA) { doTest(3, 4, CVTCODE(BGR2RGBA)); }
+OCL_TEST_P(CvtColor, RGB2BGRA) { doTest(3, 4, CVTCODE(BGR2RGBA)); }
+OCL_TEST_P(CvtColor, RGBA2BGR) { doTest(4, 3, CVTCODE(RGBA2BGR)); }
+OCL_TEST_P(CvtColor, BGRA2RGB) { doTest(4, 3, CVTCODE(RGBA2BGR)); }
+OCL_TEST_P(CvtColor, BGR2RGB) { doTest(3, 3, CVTCODE(BGR2RGB)); }
+OCL_TEST_P(CvtColor, RGB2BGR) { doTest(3, 3, CVTCODE(BGR2RGB)); }
+OCL_TEST_P(CvtColor, BGRA2RGBA) { doTest(4, 4, CVTCODE(BGRA2RGBA)); }
+OCL_TEST_P(CvtColor, RGBA2BGRA) { doTest(4, 4, CVTCODE(BGRA2RGBA)); }
 
-OCL_TEST_P(CvtColor, BGR2GRAY)
-{
-    doTest(3, 1, CVTCODE(BGR2GRAY));
-}
-OCL_TEST_P(CvtColor, GRAY2BGR)
-{
-    doTest(1, 3, CVTCODE(GRAY2BGR));
-}
-
-OCL_TEST_P(CvtColor, RGBA2GRAY)
-{
-    doTest(4, 1, CVTCODE(RGBA2GRAY));
-}
-OCL_TEST_P(CvtColor, GRAY2RGBA)
-{
-    doTest(1, 4, CVTCODE(GRAY2RGBA));
-}
+// RGB <-> Gray
 
-OCL_TEST_P(CvtColor, BGRA2GRAY)
-{
-    doTest(4, 1, CVTCODE(BGRA2GRAY));
-}
-OCL_TEST_P(CvtColor, GRAY2BGRA)
-{
-    doTest(1, 4, CVTCODE(GRAY2BGRA));
-}
+OCL_TEST_P(CvtColor, RGB2GRAY) { doTest(3, 1, CVTCODE(RGB2GRAY)); }
+OCL_TEST_P(CvtColor, GRAY2RGB) { doTest(1, 3, CVTCODE(GRAY2RGB)); }
+OCL_TEST_P(CvtColor, BGR2GRAY) { doTest(3, 1, CVTCODE(BGR2GRAY)); }
+OCL_TEST_P(CvtColor, GRAY2BGR) { doTest(1, 3, CVTCODE(GRAY2BGR)); }
+OCL_TEST_P(CvtColor, RGBA2GRAY) { doTest(4, 1, CVTCODE(RGBA2GRAY)); }
+OCL_TEST_P(CvtColor, GRAY2RGBA) { doTest(1, 4, CVTCODE(GRAY2RGBA)); }
+OCL_TEST_P(CvtColor, BGRA2GRAY) { doTest(4, 1, CVTCODE(BGRA2GRAY)); }
+OCL_TEST_P(CvtColor, GRAY2BGRA) { doTest(1, 4, CVTCODE(GRAY2BGRA)); }
 
 // RGB <-> YUV
 
-OCL_TEST_P(CvtColor, RGB2YUV)
-{
-    doTest(3, 3, CVTCODE(RGB2YUV));
-}
-OCL_TEST_P(CvtColor, BGR2YUV)
-{
-    doTest(3, 3, CVTCODE(BGR2YUV));
-}
-OCL_TEST_P(CvtColor, RGBA2YUV)
-{
-    doTest(4, 3, CVTCODE(RGB2YUV));
-}
-OCL_TEST_P(CvtColor, BGRA2YUV)
-{
-    doTest(4, 3, CVTCODE(BGR2YUV));
-}
-OCL_TEST_P(CvtColor, YUV2RGB)
-{
-    doTest(3, 3, CVTCODE(YUV2RGB));
-}
-OCL_TEST_P(CvtColor, YUV2BGR)
-{
-    doTest(3, 3, CVTCODE(YUV2BGR));
-}
-OCL_TEST_P(CvtColor, YUV2RGBA)
-{
-    doTest(3, 4, CVTCODE(YUV2RGB));
-}
-OCL_TEST_P(CvtColor, YUV2BGRA)
-{
-    doTest(3, 4, CVTCODE(YUV2BGR));
-}
+OCL_TEST_P(CvtColor, RGB2YUV) { doTest(3, 3, CVTCODE(RGB2YUV)); }
+OCL_TEST_P(CvtColor, BGR2YUV) { doTest(3, 3, CVTCODE(BGR2YUV)); }
+OCL_TEST_P(CvtColor, RGBA2YUV) { doTest(4, 3, CVTCODE(RGB2YUV)); }
+OCL_TEST_P(CvtColor, BGRA2YUV) { doTest(4, 3, CVTCODE(BGR2YUV)); }
+OCL_TEST_P(CvtColor, YUV2RGB) { doTest(3, 3, CVTCODE(YUV2RGB)); }
+OCL_TEST_P(CvtColor, YUV2BGR) { doTest(3, 3, CVTCODE(YUV2BGR)); }
+OCL_TEST_P(CvtColor, YUV2RGBA) { doTest(3, 4, CVTCODE(YUV2RGB)); }
+OCL_TEST_P(CvtColor, YUV2BGRA) { doTest(3, 4, CVTCODE(YUV2BGR)); }
 
 // RGB <-> YCrCb
 
-OCL_TEST_P(CvtColor, RGB2YCrCb)
-{
-    doTest(3, 3, CVTCODE(RGB2YCrCb));
-}
-OCL_TEST_P(CvtColor, BGR2YCrCb)
-{
-    doTest(3, 3, CVTCODE(BGR2YCrCb));
-}
-OCL_TEST_P(CvtColor, RGBA2YCrCb)
-{
-    doTest(4, 3, CVTCODE(RGB2YCrCb));
-}
-OCL_TEST_P(CvtColor, BGRA2YCrCb)
-{
-    doTest(4, 3, CVTCODE(BGR2YCrCb));
-}
-OCL_TEST_P(CvtColor, YCrCb2RGB)
-{
-    doTest(3, 3, CVTCODE(YCrCb2RGB));
-}
-OCL_TEST_P(CvtColor, YCrCb2BGR)
-{
-    doTest(3, 3, CVTCODE(YCrCb2BGR));
-}
-OCL_TEST_P(CvtColor, YCrCb2RGBA)
-{
-    doTest(3, 4, CVTCODE(YCrCb2RGB));
-}
-OCL_TEST_P(CvtColor, YCrCb2BGRA)
-{
-    doTest(3, 4, CVTCODE(YCrCb2BGR));
-}
+OCL_TEST_P(CvtColor, RGB2YCrCb) { doTest(3, 3, CVTCODE(RGB2YCrCb)); }
+OCL_TEST_P(CvtColor, BGR2YCrCb) { doTest(3, 3, CVTCODE(BGR2YCrCb)); }
+OCL_TEST_P(CvtColor, RGBA2YCrCb) { doTest(4, 3, CVTCODE(RGB2YCrCb)); }
+OCL_TEST_P(CvtColor, BGRA2YCrCb) { doTest(4, 3, CVTCODE(BGR2YCrCb)); }
+OCL_TEST_P(CvtColor, YCrCb2RGB) { doTest(3, 3, CVTCODE(YCrCb2RGB)); }
+OCL_TEST_P(CvtColor, YCrCb2BGR) { doTest(3, 3, CVTCODE(YCrCb2BGR)); }
+OCL_TEST_P(CvtColor, YCrCb2RGBA) { doTest(3, 4, CVTCODE(YCrCb2RGB)); }
+OCL_TEST_P(CvtColor, YCrCb2BGRA) { doTest(3, 4, CVTCODE(YCrCb2BGR)); }
 
 // RGB <-> XYZ
 
-OCL_TEST_P(CvtColor, RGB2XYZ)
-{
-    doTest(3, 3, CVTCODE(RGB2XYZ));
-}
-OCL_TEST_P(CvtColor, BGR2XYZ)
-{
-    doTest(3, 3, CVTCODE(BGR2XYZ));
-}
-OCL_TEST_P(CvtColor, RGBA2XYZ)
-{
-    doTest(4, 3, CVTCODE(RGB2XYZ));
-}
-OCL_TEST_P(CvtColor, BGRA2XYZ)
-{
-    doTest(4, 3, CVTCODE(BGR2XYZ));
-}
+OCL_TEST_P(CvtColor, RGB2XYZ) { doTest(3, 3, CVTCODE(RGB2XYZ)); }
+OCL_TEST_P(CvtColor, BGR2XYZ) { doTest(3, 3, CVTCODE(BGR2XYZ)); }
+OCL_TEST_P(CvtColor, RGBA2XYZ) { doTest(4, 3, CVTCODE(RGB2XYZ)); }
+OCL_TEST_P(CvtColor, BGRA2XYZ) { doTest(4, 3, CVTCODE(BGR2XYZ)); }
 
-OCL_TEST_P(CvtColor, XYZ2RGB)
-{
-    doTest(3, 3, CVTCODE(XYZ2RGB));
-}
-OCL_TEST_P(CvtColor, XYZ2BGR)
-{
-    doTest(3, 3, CVTCODE(XYZ2BGR));
-}
-OCL_TEST_P(CvtColor, XYZ2RGBA)
-{
-    doTest(3, 4, CVTCODE(XYZ2RGB));
-}
-OCL_TEST_P(CvtColor, XYZ2BGRA)
-{
-    doTest(3, 4, CVTCODE(XYZ2BGR));
-}
+OCL_TEST_P(CvtColor, XYZ2RGB) { doTest(3, 3, CVTCODE(XYZ2RGB)); }
+OCL_TEST_P(CvtColor, XYZ2BGR) { doTest(3, 3, CVTCODE(XYZ2BGR)); }
+OCL_TEST_P(CvtColor, XYZ2RGBA) { doTest(3, 4, CVTCODE(XYZ2RGB)); }
+OCL_TEST_P(CvtColor, XYZ2BGRA) { doTest(3, 4, CVTCODE(XYZ2BGR)); }
 
 // YUV -> RGBA_NV12
 
@@ -289,25 +205,10 @@ struct CvtColor_YUV420 :
     }
 };
 
-OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12)
-{
-    doTest(1, 4, CV_YUV2RGBA_NV12);
-}
-
-OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12)
-{
-    doTest(1, 4, CV_YUV2BGRA_NV12);
-}
-
-OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12)
-{
-    doTest(1, 3, CV_YUV2RGB_NV12);
-}
-
-OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12)
-{
-    doTest(1, 3, CV_YUV2BGR_NV12);
-}
+OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) { doTest(1, 4, CV_YUV2RGBA_NV12); }
+OCL_TEST_P(CvtColor_YUV420, YUV2BGRA_NV12) { doTest(1, 4, CV_YUV2BGRA_NV12); }
+OCL_TEST_P(CvtColor_YUV420, YUV2RGB_NV12) { doTest(1, 3, CV_YUV2RGB_NV12); }
+OCL_TEST_P(CvtColor_YUV420, YUV2BGR_NV12) { doTest(1, 3, CV_YUV2BGR_NV12); }
 
 INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor,
                             testing::Combine(