From 3cc9502c90836c3fcda86edfb29dbf301477ae6b Mon Sep 17 00:00:00 2001 From: Ilya Lavrenov Date: Sun, 10 Nov 2013 20:56:18 +0400 Subject: [PATCH] added RGB[A] <-> BGR[A] conversion to ocl::cvtColor --- modules/ocl/src/color.cpp | 59 ++++++++--- modules/ocl/src/opencl/cvt_color.cl | 51 +++++++-- modules/ocl/test/test_color.cpp | 201 +++++++++--------------------------- 3 files changed, 136 insertions(+), 175 deletions(-) diff --git a/modules/ocl/src/color.cpp b/modules/ocl/src/color.cpp index 9688be0..f27366a 100644 --- a/modules/ocl/src/color.cpp +++ b/modules/ocl/src/color.cpp @@ -79,7 +79,7 @@ static void fromRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std:: static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::string & kernelName, const oclMat & data = oclMat()) { - std::string build_options = format("-D DEPTH_%d -D channels=%d", src.depth(), dst.channels()); + std::string build_options = format("-D DEPTH_%d -D dcn=%d", src.depth(), dst.channels()); int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); @@ -101,6 +101,27 @@ static void toRGB_caller(const oclMat &src, oclMat &dst, int bidx, const std::st openCLExecuteKernel(src.clCxt, &cvt_color, kernelName.c_str(), gt, lt, args, -1, -1, build_options.c_str()); } +static void RGB_caller(const oclMat &src, oclMat &dst, bool reverse) +{ + std::string build_options = format("-D DEPTH_%d -D dcn=%d -D scn=%d -D %s", src.depth(), + dst.channels(), src.channels(), reverse ? "REVERSE" : "ORDER"); + int src_offset = src.offset / src.elemSize1(), src_step = src.step1(); + int dst_offset = dst.offset / dst.elemSize1(), dst_step = dst.step1(); + + 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_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, "RGB", 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(); @@ -110,18 +131,24 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) switch (code) { - /* - case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR: - case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA: - case CV_BGR2BGR565: case CV_BGR2BGR555: case CV_RGB2BGR565: case CV_RGB2BGR555: - case CV_BGRA2BGR565: case CV_BGRA2BGR555: case CV_RGBA2BGR565: case CV_RGBA2BGR555: - case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: - case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: - */ - case CV_RGB2GRAY: - case CV_BGR2GRAY: - case CV_RGBA2GRAY: - case CV_BGRA2GRAY: + case CV_BGR2BGRA: case CV_RGB2BGRA: case CV_BGRA2BGR: + case CV_RGBA2BGR: case CV_RGB2BGR: case CV_BGRA2RGBA: + { + CV_Assert(scn == 3 || scn == 4); + dcn = code == CV_BGR2BGRA || code == CV_RGB2BGRA || code == CV_BGRA2RGBA ? 4 : 3; + bool reverse = !(code == CV_BGR2BGRA || code == CV_BGRA2BGR); + dst.create(sz, CV_MAKE_TYPE(depth, 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: + case CV_BGR5652BGR: case CV_BGR5552BGR: case CV_BGR5652RGB: case CV_BGR5552RGB: + case CV_BGR5652BGRA: case CV_BGR5552BGRA: case CV_BGR5652RGBA: case CV_BGR5552RGBA: + */ + case CV_RGB2GRAY: case CV_BGR2GRAY: + case CV_RGBA2GRAY: case CV_BGRA2GRAY: { CV_Assert(scn == 3 || scn == 4); bidx = code == CV_BGR2GRAY || code == CV_BGRA2GRAY ? 0 : 2; @@ -158,10 +185,8 @@ static void cvtColor_caller(const oclMat &src, oclMat &dst, int code, int dcn) toRGB_caller(src, dst, bidx, "YUV2RGB"); break; } - case CV_YUV2RGB_NV12: - case CV_YUV2BGR_NV12: - case CV_YUV2RGBA_NV12: - case CV_YUV2BGRA_NV12: + case CV_YUV2RGB_NV12: case CV_YUV2BGR_NV12: + case CV_YUV2RGBA_NV12: case CV_YUV2BGRA_NV12: { CV_Assert(scn == 1); CV_Assert( sz.width % 2 == 0 && sz.height % 3 == 0 && depth == CV_8U ); diff --git a/modules/ocl/src/opencl/cvt_color.cl b/modules/ocl/src/opencl/cvt_color.cl index 0f44897..916d44b 100644 --- a/modules/ocl/src/opencl/cvt_color.cl +++ b/modules/ocl/src/opencl/cvt_color.cl @@ -119,8 +119,8 @@ __kernel void Gray2RGB(int cols, int rows, int src_step, int dst_step, int bidx, dst[dst_idx] = val; dst[dst_idx + 1] = val; dst[dst_idx + 2] = val; -#if channels == 4 - dst[dst_idx + 3] = MAX_NUM; +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; #endif } } @@ -195,8 +195,8 @@ __kernel void YUV2RGB(int cols, int rows, int src_step, int dst_step, dst[dst_idx + bidx] = SAT_CAST( b ); dst[dst_idx + 1] = SAT_CAST( g ); dst[dst_idx + (bidx^2)] = SAT_CAST( r ); -#if channels == 4 - dst[dst_idx + 3] = MAX_NUM; +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; #endif } } @@ -332,8 +332,8 @@ __kernel void YCrCb2RGB(int cols, int rows, int src_step, int dst_step, dst[dst_idx + (bidx^2)] = SAT_CAST(r); dst[dst_idx + 1] = SAT_CAST(g); dst[dst_idx + bidx] = SAT_CAST(b); -#if channels == 4 - dst[dst_idx + 3] = MAX_NUM; +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; #endif } } @@ -397,8 +397,43 @@ __kernel void XYZ2RGB(int cols, int rows, int src_step, int dst_step, dst[dst_idx] = SAT_CAST(b); dst[dst_idx + 1] = SAT_CAST(g); dst[dst_idx + 2] = SAT_CAST(r); -#if channels == 4 - dst[dst_idx + 3] = MAX_NUM; +#if dcn == 4 + dst[dst_idx + 3] = MAX_NUM; +#endif + } +} + +///////////////////////////////////// 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 } } diff --git a/modules/ocl/test/test_color.cpp b/modules/ocl/test/test_color.cpp index 977fd20..e982684 100644 --- a/modules/ocl/test/test_color.cpp +++ b/modules/ocl/test/test_color.cpp @@ -121,149 +121,65 @@ PARAM_TEST_CASE(CvtColor, MatDepth, bool) #define CVTCODE(name) COLOR_ ## name -// RGB <-> Gray - -OCL_TEST_P(CvtColor, RGB2GRAY) -{ - doTest(3, 1, CVTCODE(RGB2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2RGB) -{ - doTest(1, 3, CVTCODE(GRAY2RGB)); -} +// RGB[A] <-> BGR[A] + +OCL_TEST_P(CvtColor, BGR2BGRA) { doTest(3, 4, CVTCODE(BGR2BGRA)); } +OCL_TEST_P(CvtColor, RGB2RGBA) { doTest(3, 4, CVTCODE(BGR2BGRA)); } +OCL_TEST_P(CvtColor, BGRA2BGR) { doTest(4, 3, CVTCODE(BGRA2BGR)); } +OCL_TEST_P(CvtColor, RGBA2RGB) { doTest(4, 3, CVTCODE(BGRA2BGR)); } +OCL_TEST_P(CvtColor, BGR2RGBA) { doTest(3, 4, CVTCODE(BGR2RGBA)); } +OCL_TEST_P(CvtColor, RGB2BGRA) { doTest(3, 4, CVTCODE(BGR2RGBA)); } +OCL_TEST_P(CvtColor, RGBA2BGR) { doTest(4, 3, CVTCODE(RGBA2BGR)); } +OCL_TEST_P(CvtColor, BGRA2RGB) { doTest(4, 3, CVTCODE(RGBA2BGR)); } +OCL_TEST_P(CvtColor, BGR2RGB) { doTest(3, 3, CVTCODE(BGR2RGB)); } +OCL_TEST_P(CvtColor, RGB2BGR) { doTest(3, 3, CVTCODE(BGR2RGB)); } +OCL_TEST_P(CvtColor, BGRA2RGBA) { doTest(4, 4, CVTCODE(BGRA2RGBA)); } +OCL_TEST_P(CvtColor, RGBA2BGRA) { doTest(4, 4, CVTCODE(BGRA2RGBA)); } -OCL_TEST_P(CvtColor, BGR2GRAY) -{ - doTest(3, 1, CVTCODE(BGR2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2BGR) -{ - doTest(1, 3, CVTCODE(GRAY2BGR)); -} - -OCL_TEST_P(CvtColor, RGBA2GRAY) -{ - doTest(4, 1, CVTCODE(RGBA2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2RGBA) -{ - doTest(1, 4, CVTCODE(GRAY2RGBA)); -} +// RGB <-> Gray -OCL_TEST_P(CvtColor, BGRA2GRAY) -{ - doTest(4, 1, CVTCODE(BGRA2GRAY)); -} -OCL_TEST_P(CvtColor, GRAY2BGRA) -{ - doTest(1, 4, CVTCODE(GRAY2BGRA)); -} +OCL_TEST_P(CvtColor, RGB2GRAY) { doTest(3, 1, CVTCODE(RGB2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2RGB) { doTest(1, 3, CVTCODE(GRAY2RGB)); } +OCL_TEST_P(CvtColor, BGR2GRAY) { doTest(3, 1, CVTCODE(BGR2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2BGR) { doTest(1, 3, CVTCODE(GRAY2BGR)); } +OCL_TEST_P(CvtColor, RGBA2GRAY) { doTest(4, 1, CVTCODE(RGBA2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2RGBA) { doTest(1, 4, CVTCODE(GRAY2RGBA)); } +OCL_TEST_P(CvtColor, BGRA2GRAY) { doTest(4, 1, CVTCODE(BGRA2GRAY)); } +OCL_TEST_P(CvtColor, GRAY2BGRA) { doTest(1, 4, CVTCODE(GRAY2BGRA)); } // RGB <-> YUV -OCL_TEST_P(CvtColor, RGB2YUV) -{ - doTest(3, 3, CVTCODE(RGB2YUV)); -} -OCL_TEST_P(CvtColor, BGR2YUV) -{ - doTest(3, 3, CVTCODE(BGR2YUV)); -} -OCL_TEST_P(CvtColor, RGBA2YUV) -{ - doTest(4, 3, CVTCODE(RGB2YUV)); -} -OCL_TEST_P(CvtColor, BGRA2YUV) -{ - doTest(4, 3, CVTCODE(BGR2YUV)); -} -OCL_TEST_P(CvtColor, YUV2RGB) -{ - doTest(3, 3, CVTCODE(YUV2RGB)); -} -OCL_TEST_P(CvtColor, YUV2BGR) -{ - doTest(3, 3, CVTCODE(YUV2BGR)); -} -OCL_TEST_P(CvtColor, YUV2RGBA) -{ - doTest(3, 4, CVTCODE(YUV2RGB)); -} -OCL_TEST_P(CvtColor, YUV2BGRA) -{ - doTest(3, 4, CVTCODE(YUV2BGR)); -} +OCL_TEST_P(CvtColor, RGB2YUV) { doTest(3, 3, CVTCODE(RGB2YUV)); } +OCL_TEST_P(CvtColor, BGR2YUV) { doTest(3, 3, CVTCODE(BGR2YUV)); } +OCL_TEST_P(CvtColor, RGBA2YUV) { doTest(4, 3, CVTCODE(RGB2YUV)); } +OCL_TEST_P(CvtColor, BGRA2YUV) { doTest(4, 3, CVTCODE(BGR2YUV)); } +OCL_TEST_P(CvtColor, YUV2RGB) { doTest(3, 3, CVTCODE(YUV2RGB)); } +OCL_TEST_P(CvtColor, YUV2BGR) { doTest(3, 3, CVTCODE(YUV2BGR)); } +OCL_TEST_P(CvtColor, YUV2RGBA) { doTest(3, 4, CVTCODE(YUV2RGB)); } +OCL_TEST_P(CvtColor, YUV2BGRA) { doTest(3, 4, CVTCODE(YUV2BGR)); } // RGB <-> YCrCb -OCL_TEST_P(CvtColor, RGB2YCrCb) -{ - doTest(3, 3, CVTCODE(RGB2YCrCb)); -} -OCL_TEST_P(CvtColor, BGR2YCrCb) -{ - doTest(3, 3, CVTCODE(BGR2YCrCb)); -} -OCL_TEST_P(CvtColor, RGBA2YCrCb) -{ - doTest(4, 3, CVTCODE(RGB2YCrCb)); -} -OCL_TEST_P(CvtColor, BGRA2YCrCb) -{ - doTest(4, 3, CVTCODE(BGR2YCrCb)); -} -OCL_TEST_P(CvtColor, YCrCb2RGB) -{ - doTest(3, 3, CVTCODE(YCrCb2RGB)); -} -OCL_TEST_P(CvtColor, YCrCb2BGR) -{ - doTest(3, 3, CVTCODE(YCrCb2BGR)); -} -OCL_TEST_P(CvtColor, YCrCb2RGBA) -{ - doTest(3, 4, CVTCODE(YCrCb2RGB)); -} -OCL_TEST_P(CvtColor, YCrCb2BGRA) -{ - doTest(3, 4, CVTCODE(YCrCb2BGR)); -} +OCL_TEST_P(CvtColor, RGB2YCrCb) { doTest(3, 3, CVTCODE(RGB2YCrCb)); } +OCL_TEST_P(CvtColor, BGR2YCrCb) { doTest(3, 3, CVTCODE(BGR2YCrCb)); } +OCL_TEST_P(CvtColor, RGBA2YCrCb) { doTest(4, 3, CVTCODE(RGB2YCrCb)); } +OCL_TEST_P(CvtColor, BGRA2YCrCb) { doTest(4, 3, CVTCODE(BGR2YCrCb)); } +OCL_TEST_P(CvtColor, YCrCb2RGB) { doTest(3, 3, CVTCODE(YCrCb2RGB)); } +OCL_TEST_P(CvtColor, YCrCb2BGR) { doTest(3, 3, CVTCODE(YCrCb2BGR)); } +OCL_TEST_P(CvtColor, YCrCb2RGBA) { doTest(3, 4, CVTCODE(YCrCb2RGB)); } +OCL_TEST_P(CvtColor, YCrCb2BGRA) { doTest(3, 4, CVTCODE(YCrCb2BGR)); } // RGB <-> XYZ -OCL_TEST_P(CvtColor, RGB2XYZ) -{ - doTest(3, 3, CVTCODE(RGB2XYZ)); -} -OCL_TEST_P(CvtColor, BGR2XYZ) -{ - doTest(3, 3, CVTCODE(BGR2XYZ)); -} -OCL_TEST_P(CvtColor, RGBA2XYZ) -{ - doTest(4, 3, CVTCODE(RGB2XYZ)); -} -OCL_TEST_P(CvtColor, BGRA2XYZ) -{ - doTest(4, 3, CVTCODE(BGR2XYZ)); -} +OCL_TEST_P(CvtColor, RGB2XYZ) { doTest(3, 3, CVTCODE(RGB2XYZ)); } +OCL_TEST_P(CvtColor, BGR2XYZ) { doTest(3, 3, CVTCODE(BGR2XYZ)); } +OCL_TEST_P(CvtColor, RGBA2XYZ) { doTest(4, 3, CVTCODE(RGB2XYZ)); } +OCL_TEST_P(CvtColor, BGRA2XYZ) { doTest(4, 3, CVTCODE(BGR2XYZ)); } -OCL_TEST_P(CvtColor, XYZ2RGB) -{ - doTest(3, 3, CVTCODE(XYZ2RGB)); -} -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)); -} +OCL_TEST_P(CvtColor, XYZ2RGB) { doTest(3, 3, CVTCODE(XYZ2RGB)); } +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)); } // YUV -> RGBA_NV12 @@ -289,25 +205,10 @@ struct CvtColor_YUV420 : } }; -OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) -{ - doTest(1, 4, CV_YUV2RGBA_NV12); -} - -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); -} +OCL_TEST_P(CvtColor_YUV420, YUV2RGBA_NV12) { doTest(1, 4, CV_YUV2RGBA_NV12); } +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, CvtColor, testing::Combine( -- 2.7.4