diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index f27366a6ae..f9bccc89b6 100644 --- a/modules/ocl/src/color.cpp +++ b/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 > 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 > 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: { diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 916d44bf93..b0bd7d0848 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/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 + } +} diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index e98268425d..105ee89d90 100644 --- a/modules/ocl/test/test_color.cpp +++ b/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)),