added RGB5x5 <-> RGB conversion

modules/ocl/src/color.cpp

@@ -122,6 +122,50 @@ static void RGB_caller(const oclMat &src, oclMat &dst, bool reverse)
     openCLExecuteKernel(src.clCxt, &cvt_color, "RGB", gt, lt, args, -1, -1, build_options.c_str());
 }
 
+static void RGB5x52RGB_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits)
+{
+    std::string build_options = format("-D DEPTH_%d -D greenbits=%d -D dcn=%d",
+                                       src.depth(), greenbits, dst.channels());
+    int src_offset = src.offset >> 1, src_step = src.step >> 1;
+    int dst_offset = (int)dst.offset, dst_step = (int)dst.step;
+
+    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_int) , (void *)&bidx));
+    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, "RGB5x52RGB", gt, lt, args, -1, -1, build_options.c_str());
+}
+
+static void RGB2RGB5x5_caller(const oclMat &src, oclMat &dst, int bidx, int greenbits)
+{
+    std::string build_options = format("-D DEPTH_%d -D greenbits=%d -D scn=%d",
+                                       src.depth(), greenbits, src.channels());
+    int src_offset = (int)src.offset, src_step = (int)src.step;
+    int dst_offset = dst.offset >> 1, dst_step = dst.step >> 1;
+
+    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_int) , (void *)&bidx));
+    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, "RGB2RGB5x5", 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();
@@ -141,12 +185,31 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int 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:
+    {
+        CV_Assert((scn == 3 || scn == 4) && depth == CV_8U );
+        bidx = code == CV_BGR2BGR565 || code == CV_BGR2BGR555 ||
+            code == CV_BGRA2BGR565 || code == CV_BGRA2BGR555 ? 0 : 2;
+        int greenbits = code == CV_BGR2BGR565 || code == CV_RGB2BGR565 ||
+            code == CV_BGRA2BGR565 || code == CV_RGBA2BGR565 ? 6 : 5;
+        dst.create(sz, CV_8UC2);
+        RGB2RGB5x5_caller(src, dst, bidx, greenbits);
+        break;
+    }
     case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB:
     case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA:
-    */
+    {
+        dcn = code == CV_BGR5652BGRA || code == CV_BGR5552BGRA || code == CV_BGR5652RGBA || code == CV_BGR5552RGBA ? 4 : 3;
+        CV_Assert((dcn == 3 || dcn == 4) && scn == 2 && depth == CV_8U);
+        bidx = code == CV_BGR5652BGR || code == CV_BGR5552BGR ||
+            code == CV_BGR5652BGRA || code == CV_BGR5552BGRA ? 0 : 2;
+        int greenbits = code == CV_BGR5652BGR || code == CV_BGR5652RGB ||
+            code == CV_BGR5652BGRA || code == CV_BGR5652RGBA ? 6 : 5;
+        dst.create(sz, CV_MAKETYPE(depth, dcn));
+        RGB5x52RGB_caller(src, dst, bidx, greenbits);
+        break;
+    }
     case CV_RGB2GRAY: case CV_BGR2GRAY:
     case CV_RGBA2GRAY: case CV_BGRA2GRAY:
     {

modules/ocl/src/opencl/cvt_color.cl

@@ -437,3 +437,61 @@ __kernel void RGB(int cols, int rows, int src_step, int dst_step,
 #endif
     }
 }
+
+///////////////////////////////////// RGB5x5 <-> RGB //////////////////////////////////////
+
+__kernel void RGB5x52RGB(int cols, int rows, int src_step, int dst_step, int bidx,
+                         __global const ushort * src, __global uchar * 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);
+        int dst_idx = mad24(y, dst_step, dst_offset + (x << 2));
+        ushort t = 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/ocl/test/test_color.cpp

@@ -181,6 +181,30 @@ 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)); }
 
+// RGB5x5 <-> RGB
+
+typedef CvtColor CvtColor8u;
+
+OCL_TEST_P(CvtColor8u, BGR5652BGR) { doTest(2, 3, CVTCODE(BGR5652BGR)); }
+OCL_TEST_P(CvtColor8u, BGR5652RGB) { doTest(2, 3, CVTCODE(BGR5652RGB)); }
+OCL_TEST_P(CvtColor8u, BGR5652BGRA) { doTest(2, 4, CVTCODE(BGR5652BGRA)); }
+OCL_TEST_P(CvtColor8u, BGR5652RGBA) { doTest(2, 4, CVTCODE(BGR5652RGBA)); }
+
+OCL_TEST_P(CvtColor8u, BGR5552BGR) { doTest(2, 3, CVTCODE(BGR5552BGR)); }
+OCL_TEST_P(CvtColor8u, BGR5552RGB) { doTest(2, 3, CVTCODE(BGR5552RGB)); }
+OCL_TEST_P(CvtColor8u, BGR5552BGRA) { doTest(2, 4, CVTCODE(BGR5552BGRA)); }
+OCL_TEST_P(CvtColor8u, BGR5552RGBA) { doTest(2, 4, CVTCODE(BGR5552RGBA)); }
+
+OCL_TEST_P(CvtColor8u, BGR2BGR565) { doTest(3, 2, CVTCODE(BGR2BGR565)); }
+OCL_TEST_P(CvtColor8u, RGB2BGR565) { doTest(3, 2, CVTCODE(RGB2BGR565)); }
+OCL_TEST_P(CvtColor8u, BGRA2BGR565) { doTest(4, 2, CVTCODE(BGRA2BGR565)); }
+OCL_TEST_P(CvtColor8u, RGBA2BGR565) { doTest(4, 2, CVTCODE(RGBA2BGR565)); }
+
+OCL_TEST_P(CvtColor8u, BGR2BGR555) { doTest(3, 2, CVTCODE(BGR2BGR555)); }
+OCL_TEST_P(CvtColor8u, RGB2BGR555) { doTest(3, 2, CVTCODE(RGB2BGR555)); }
+OCL_TEST_P(CvtColor8u, BGRA2BGR555) { doTest(4, 2, CVTCODE(BGRA2BGR555)); }
+OCL_TEST_P(CvtColor8u, RGBA2BGR555) { doTest(4, 2, CVTCODE(RGBA2BGR555)); }
+
 // YUV -> RGBA_NV12
 
 struct CvtColor_YUV420 :
@@ -210,6 +234,10 @@ 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, CvtColor8u,
+                            testing::Combine(testing::Values(MatDepth(CV_8U)), Bool()));
+
 INSTANTIATE_TEST_CASE_P(OCL_ImgProc, CvtColor,
                             testing::Combine(
                                 testing::Values(MatDepth(CV_8U), MatDepth(CV_16U), MatDepth(CV_32F)),