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();
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:
{
#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
+ }
+}
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 :
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)),