RGB[A] -> RGB5x5

modules/imgproc/src/color.cpp

@@ -2706,8 +2706,6 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
     /*
      case COLOR_BGR2BGR565: case COLOR_BGR2BGR555: case COLOR_RGB2BGR565: case COLOR_RGB2BGR555:
      case COLOR_BGRA2BGR565: case COLOR_BGRA2BGR555: case COLOR_RGBA2BGR565: case COLOR_RGBA2BGR555:
-     case COLOR_BGR5652BGR: case COLOR_BGR5552BGR: case COLOR_BGR5652RGB: case COLOR_BGR5552RGB:
-     case COLOR_BGR5652BGRA: case COLOR_BGR5552BGRA: case COLOR_BGR5652RGBA: case COLOR_BGR5552RGBA:
      */
     case COLOR_BGR2BGRA: case COLOR_RGB2BGRA: case COLOR_BGRA2BGR:
     case COLOR_RGBA2BGR: case COLOR_RGB2BGR: case COLOR_BGRA2RGBA:
@@ -2720,6 +2718,19 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
                         reverse ? "REVERSE" : "ORDER"));
         break;
     }
+    case COLOR_BGR5652BGR: case COLOR_BGR5552BGR: case COLOR_BGR5652RGB: case COLOR_BGR5552RGB:
+    case COLOR_BGR5652BGRA: case COLOR_BGR5552BGRA: case COLOR_BGR5652RGBA: case COLOR_BGR5552RGBA:
+    {
+        dcn = code == COLOR_BGR5652BGRA || code == COLOR_BGR5552BGRA || code == COLOR_BGR5652RGBA || code == COLOR_BGR5552RGBA ? 4 : 3;
+        CV_Assert((dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U);
+        bidx = code == COLOR_BGR5652BGR || code == COLOR_BGR5552BGR ||
+            code == COLOR_BGR5652BGRA || code == COLOR_BGR5552BGRA ? 0 : 2;
+        int greenbits = code == COLOR_BGR5652BGR || code == COLOR_BGR5652RGB ||
+            code == COLOR_BGR5652BGRA || code == COLOR_BGR5652RGBA ? 6 : 5;
+        k.create("RGB5x52RGB", ocl::imgproc::cvtcolor_oclsrc,
+                 format("-D depth=%d -D scn=2 -D dcn=%s -D bidx=%d -D greenbits=%d", depth, scn, dcn, bidx, greenbits));
+        break;
+    }
     case COLOR_BGR2GRAY: case COLOR_BGRA2GRAY:
     case COLOR_RGB2GRAY: case COLOR_RGBA2GRAY:
     {

modules/imgproc/src/opencl/cvtcolor.cl

@@ -455,6 +455,63 @@ __kernel void RGB(__global const uchar* srcptr, int src_step, int src_offset,
 }
 
 
+///////////////////////////////////// RGB5x5 <-> RGB //////////////////////////////////////
+
+__kernel void RGB5x52RGB(__global const uchar* src, int src_step, int src_offset,
+                         __global uchar* dst, int dst_step, int dst_offset,
+                         int rows, int cols)
+{
+    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);
+        ushort t = *((__global const ushort*)(src + src_idx));
+
+#if greenbits == 6
+        dst[dst_idx + bidx] = (uchar)(t << 3);
+        dst[dst_idx + 1] = (uchar)((t >> 3) & ~3);
+        dst[dst_idx + (bidx^2)] = (uchar)((t >> 8) & ~7);
+#else
+        dst[dst_idx + bidx] = (uchar)(t << 3);
+        dst[dst_idx + 1] = (uchar)((t >> 2) & ~7);
+        dst[dst_idx + (bidx^2)] = (uchar)((t >> 7) & ~7);
+#endif
+
+#if dcn == 4
+#if greenbits == 6
+        dst[dst_idx + 3] = 255;
+#else
+        dst[dst_idx + 3] = t & 0x8000 ? 255 : 0;
+#endif
+#endif
+    }
+}
+
+//__kernel void RGB2RGB5x5(int cols, int rows, int src_step, int dst_step, int bidx,
+//                         __global const uchar * src, __global ushort * dst,
+//                         int src_offset, int dst_offset)
+//{
+//    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 << 2));
+//        int dst_idx = mad24(y, dst_step, dst_offset + x);
+
+//#if greenbits == 6
+//            dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~3) << 3)|((src[src_idx + (bidx^2)]&~7) << 8));
+//#elif scn == 3
+//            dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)|((src[src_idx + (bidx^2)]&~7) << 7));
+//#else
+//            dst[dst_idx] = (ushort)((src[src_idx + bidx] >> 3)|((src[src_idx + 1]&~7) << 2)|
+//                ((src[src_idx + (bidx^2)]&~7) << 7)|(src[src_idx + 3] ? 0x8000 : 0));
+//#endif
+//    }
+//}
 
 
 

modules/imgproc/test/ocl/test_color.cpp

@@ -215,17 +215,17 @@ OCL_TEST_P(CvtColor, XYZ2BGRA) { performTest(3, 4, CVTCODE(XYZ2BGR)); }
 
 // RGB5x5 <-> RGB
 
-//typedef CvtColor CvtColor8u;
+typedef CvtColor CvtColor8u;
 
-//OCL_TEST_P(CvtColor8u, BGR5652BGR) { performTest(2, 3, CVTCODE(BGR5652BGR)); }
-//OCL_TEST_P(CvtColor8u, BGR5652RGB) { performTest(2, 3, CVTCODE(BGR5652RGB)); }
-//OCL_TEST_P(CvtColor8u, BGR5652BGRA) { performTest(2, 4, CVTCODE(BGR5652BGRA)); }
-//OCL_TEST_P(CvtColor8u, BGR5652RGBA) { performTest(2, 4, CVTCODE(BGR5652RGBA)); }
+OCL_TEST_P(CvtColor8u, BGR5652BGR) { performTest(2, 3, CVTCODE(BGR5652BGR)); }
+OCL_TEST_P(CvtColor8u, BGR5652RGB) { performTest(2, 3, CVTCODE(BGR5652RGB)); }
+OCL_TEST_P(CvtColor8u, BGR5652BGRA) { performTest(2, 4, CVTCODE(BGR5652BGRA)); }
+OCL_TEST_P(CvtColor8u, BGR5652RGBA) { performTest(2, 4, CVTCODE(BGR5652RGBA)); }
 
-//OCL_TEST_P(CvtColor8u, BGR5552BGR) { performTest(2, 3, CVTCODE(BGR5552BGR)); }
-//OCL_TEST_P(CvtColor8u, BGR5552RGB) { performTest(2, 3, CVTCODE(BGR5552RGB)); }
-//OCL_TEST_P(CvtColor8u, BGR5552BGRA) { performTest(2, 4, CVTCODE(BGR5552BGRA)); }
-//OCL_TEST_P(CvtColor8u, BGR5552RGBA) { performTest(2, 4, CVTCODE(BGR5552RGBA)); }
+OCL_TEST_P(CvtColor8u, BGR5552BGR) { performTest(2, 3, CVTCODE(BGR5552BGR)); }
+OCL_TEST_P(CvtColor8u, BGR5552RGB) { performTest(2, 3, CVTCODE(BGR5552RGB)); }
+OCL_TEST_P(CvtColor8u, BGR5552BGRA) { performTest(2, 4, CVTCODE(BGR5552BGRA)); }
+OCL_TEST_P(CvtColor8u, BGR5552RGBA) { performTest(2, 4, CVTCODE(BGR5552RGBA)); }
 
 //OCL_TEST_P(CvtColor8u, BGR2BGR565) { performTest(3, 2, CVTCODE(BGR2BGR565)); }
 //OCL_TEST_P(CvtColor8u, RGB2BGR565) { performTest(3, 2, CVTCODE(RGB2BGR565)); }

modules/ocl/src/opencl/cvt_color.cl

@@ -91,40 +91,6 @@ enum
     BLOCK_SIZE = 256
 };
 
-///////////////////////////////////// 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
-    }
-}
 
 ///////////////////////////////////// RGB5x5 <-> RGB //////////////////////////////////////