From 727a5e6df471111fb0627504371b5f9591800b5e Mon Sep 17 00:00:00 2001
From: Ilya Lavrenov <ilya.lavrenov@itseez.com>
Date: Wed, 27 Nov 2013 22:19:44 +0400
Subject: [PATCH] BGR5x5 <-> Gray

---
 modules/imgproc/src/color.cpp           | 22 +++++++++---
 modules/imgproc/src/opencl/cvtcolor.cl  | 46 +++++++++++++++++++++++++
 modules/imgproc/test/ocl/test_color.cpp |  8 ++---
 3 files changed, 68 insertions(+), 8 deletions(-)

diff --git a/modules/imgproc/src/color.cpp b/modules/imgproc/src/color.cpp
index 88655757c7..f5f8118f4f 100644
--- a/modules/imgproc/src/color.cpp
+++ b/modules/imgproc/src/color.cpp
@@ -2704,10 +2704,6 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
 
     switch (code)
     {
-    /*
-     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_BGR2BGRA: case COLOR_RGB2BGRA: case COLOR_BGRA2BGR:
     case COLOR_RGBA2BGR: case COLOR_RGB2BGR: case COLOR_BGRA2RGBA:
     {
@@ -2745,6 +2741,24 @@ static bool ocl_cvtColor( InputArray _src, OutputArray _dst, int code, int dcn )
                  format("-D depth=%d -D scn=%d -D dcn=2 -D bidx=%d -D greenbits=%d", depth, scn, bidx, greenbits));
         break;
     }
+    case COLOR_BGR5652GRAY: case COLOR_BGR5552GRAY:
+    {
+        CV_Assert(scn == 2 && depth == CV_8U);
+        dcn = 1;
+        int greenbits = code == COLOR_BGR5652GRAY ? 6 : 5;
+        k.create("BGR5x52Gray", ocl::imgproc::cvtcolor_oclsrc,
+                 format("-D depth=%d -D scn=2 -D dcn=1 -D bidx=0 -D greenbits=%d", depth, greenbits));
+        break;
+    }
+    case COLOR_GRAY2BGR565: case COLOR_GRAY2BGR555:
+    {
+        CV_Assert(scn == 1 && depth == CV_8U);
+        dcn = 2;
+        int greenbits = code == COLOR_GRAY2BGR565 ? 6 : 5;
+        k.create("Gray2BGR5x5", ocl::imgproc::cvtcolor_oclsrc,
+                 format("-D depth=%d -D scn=1 -D dcn=2 -D bidx=0 -D greenbits=%d", depth, greenbits));
+        break;
+    }
     case COLOR_BGR2GRAY: case COLOR_BGRA2GRAY:
     case COLOR_RGB2GRAY: case COLOR_RGBA2GRAY:
     {
diff --git a/modules/imgproc/src/opencl/cvtcolor.cl b/modules/imgproc/src/opencl/cvtcolor.cl
index 0d8272fae0..846574160d 100644
--- a/modules/imgproc/src/opencl/cvtcolor.cl
+++ b/modules/imgproc/src/opencl/cvtcolor.cl
@@ -513,8 +513,54 @@ __kernel void RGB2RGB5x5(__global const uchar* src, int src_step, int src_offset
     }
 }
 
+///////////////////////////////////// RGB5x5 <-> Gray //////////////////////////////////////
 
+__kernel void BGR5x52Gray(__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);
+        int t = *((__global const ushort*)(src + src_idx));
+
+#if greenbits == 6
+        dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y +
+                                         ((t >> 3) & 0xfc)*G2Y +
+                                         ((t >> 8) & 0xf8)*R2Y, yuv_shift);
+#else
+        dst[dst_idx] = (uchar)CV_DESCALE(((t << 3) & 0xf8)*B2Y +
+                                         ((t >> 2) & 0xf8)*G2Y +
+                                         ((t >> 7) & 0xf8)*R2Y, yuv_shift);
+#endif
+    }
+}
 
+__kernel void Gray2BGR5x5(__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);
+        int dst_idx = mad24(y, dst_step, dst_offset + x * dcnbytes);
+        int t = src[src_idx];
+
+#if greenbits == 6
+        *((__global ushort*)(dst + dst_idx)) = (ushort)((t >> 3) | ((t & ~3) << 3) | ((t & ~7) << 8));
+#else
+        t >>= 3;
+        *((__global ushort*)(dst + dst_idx)) = (ushort)(t|(t << 5)|(t << 10));
+#endif
+    }
+}
 
 
 
diff --git a/modules/imgproc/test/ocl/test_color.cpp b/modules/imgproc/test/ocl/test_color.cpp
index 695d3d1649..0dfcae84ae 100644
--- a/modules/imgproc/test/ocl/test_color.cpp
+++ b/modules/imgproc/test/ocl/test_color.cpp
@@ -239,11 +239,11 @@ OCL_TEST_P(CvtColor8u, RGBA2BGR555) { performTest(4, 2, CVTCODE(RGBA2BGR555)); }
 
 // RGB5x5 <-> Gray
 
-//OCL_TEST_P(CvtColor8u, BGR5652GRAY) { performTest(2, 1, CVTCODE(BGR5652GRAY)); }
-//OCL_TEST_P(CvtColor8u, BGR5552GRAY) { performTest(2, 1, CVTCODE(BGR5552GRAY)); }
+OCL_TEST_P(CvtColor8u, BGR5652GRAY) { performTest(2, 1, CVTCODE(BGR5652GRAY)); }
+OCL_TEST_P(CvtColor8u, BGR5552GRAY) { performTest(2, 1, CVTCODE(BGR5552GRAY)); }
 
-//OCL_TEST_P(CvtColor8u, GRAY2BGR565) { performTest(1, 2, CVTCODE(GRAY2BGR565)); }
-//OCL_TEST_P(CvtColor8u, GRAY2BGR555) { performTest(1, 2, CVTCODE(GRAY2BGR555)); }
+OCL_TEST_P(CvtColor8u, GRAY2BGR565) { performTest(1, 2, CVTCODE(GRAY2BGR565)); }
+OCL_TEST_P(CvtColor8u, GRAY2BGR555) { performTest(1, 2, CVTCODE(GRAY2BGR555)); }
 
 // RGBA <-> mRGBA